diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h index 822ff97e26..a330ab2447 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/buffer_tracing.h @@ -122,20 +122,32 @@ typedef struct /** * @brief ROCProfiler Buffer Kernel Dispatch Tracer Record. */ -typedef struct +typedef struct rocprofiler_buffer_tracing_kernel_dispatch_record_t { - uint64_t size; ///< size of this struct - rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH - rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record - rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds - rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds - rocprofiler_agent_id_t agent_id; ///< agent kernel was dispatched on - rocprofiler_queue_id_t queue_id; ///< queue kernel was dispatched on - rocprofiler_kernel_id_t kernel_id; ///< identifier for kernel - uint32_t private_segment_size; /// runtime private memory segment size - uint32_t group_segment_size; /// runtime group memory segment size - rocprofiler_dim3_t workgroup_size; /// runtime workgroup size (grid * threads) - rocprofiler_dim3_t grid_size; /// runtime grid size + uint64_t size; ///< size of this struct + rocprofiler_buffer_tracing_kind_t kind; ///< ::ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH + rocprofiler_kernel_dispatch_operation_t operation; ///< + rocprofiler_correlation_id_t correlation_id; ///< correlation ids for record + rocprofiler_thread_id_t thread_id; ///< id for thread that launched kernel + rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds + rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds + rocprofiler_agent_id_t agent_id; ///< agent kernel was dispatched on + rocprofiler_queue_id_t queue_id; ///< queue kernel was dispatched on + rocprofiler_kernel_id_t kernel_id; ///< identifier for kernel + rocprofiler_dispatch_id_t dispatch_id; ///< unique id for each dispatch + uint32_t private_segment_size; + uint32_t group_segment_size; + rocprofiler_dim3_t workgroup_size; + rocprofiler_dim3_t grid_size; + + /// @var private_segment_size + /// @brief runtime private memory segment size + /// @var group_segment_size + /// @brief runtime group memory segment size + /// @var workgroup_size + /// @brief runtime workgroup size (grid * threads) + /// @var grid_size + /// @brief runtime grid size } rocprofiler_buffer_tracing_kernel_dispatch_record_t; /** diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/callback_tracing.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/callback_tracing.h index 9dd51d62bb..eefb353d74 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/callback_tracing.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/callback_tracing.h @@ -156,10 +156,38 @@ typedef struct } rocprofiler_callback_tracing_code_object_kernel_symbol_register_data_t; +/** + * @brief ROCProfiler Kernel Dispatch Callback Tracer Record. + * + */ +typedef struct rocprofiler_callback_tracing_kernel_dispatch_data_t +{ + uint64_t size; ///< size of this struct + rocprofiler_timestamp_t start_timestamp; ///< start time in nanoseconds + rocprofiler_timestamp_t end_timestamp; ///< end time in nanoseconds + rocprofiler_agent_id_t agent_id; ///< agent kernel was dispatched on + rocprofiler_queue_id_t queue_id; ///< queue kernel was dispatched on + rocprofiler_kernel_id_t kernel_id; ///< identifier for kernel + rocprofiler_dispatch_id_t dispatch_id; ///< unique id for each dispatch + uint32_t private_segment_size; + uint32_t group_segment_size; + rocprofiler_dim3_t workgroup_size; + rocprofiler_dim3_t grid_size; + + /// @var private_segment_size + /// @brief runtime private memory segment size + /// @var group_segment_size + /// @brief runtime group memory segment size + /// @var workgroup_size + /// @brief runtime workgroup size (grid * threads) + /// @var grid_size + /// @brief runtime grid size +} rocprofiler_callback_tracing_kernel_dispatch_data_t; + /** * @brief ROCProfiler Scratch Memory Callback Data. */ -typedef struct +typedef struct rocprofiler_callback_tracing_scratch_memory_data_t { uint64_t size; ///< size of this struct rocprofiler_agent_id_t agent_id; diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h index c85dac8966..b6aa204900 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/dispatch_profile.h @@ -44,14 +44,15 @@ ROCPROFILER_EXTERN_C_INIT typedef struct rocprofiler_profile_counting_dispatch_data_t { uint64_t size; ///< Size of this struct - rocprofiler_kernel_id_t kernel_id; ///< Kernel identifier rocprofiler_agent_id_t agent_id; ///< Agent ID where kernel is launched rocprofiler_queue_id_t queue_id; ///< Queue ID where kernel packet is enqueued + rocprofiler_kernel_id_t kernel_id; ///< Kernel identifier + rocprofiler_dispatch_id_t dispatch_id; ///< unique id for each dispatch rocprofiler_correlation_id_t correlation_id; ///< Correlation ID for this dispatch - uint32_t private_segment_size; /// runtime private memory segment size - uint32_t group_segment_size; /// runtime group memory segment size - rocprofiler_dim3_t workgroup_size; /// runtime workgroup size (grid * threads) - rocprofiler_dim3_t grid_size; /// runtime grid size + uint32_t private_segment_size; ///< runtime private memory segment size + uint32_t group_segment_size; ///< runtime group memory segment size + rocprofiler_dim3_t workgroup_size; ///< runtime workgroup size (grid * threads) + rocprofiler_dim3_t grid_size; ///< runtime grid size } rocprofiler_profile_counting_dispatch_data_t; /** diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/external_correlation.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/external_correlation.h index c8391e4f42..86bede2675 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/external_correlation.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/external_correlation.h @@ -34,6 +34,99 @@ ROCPROFILER_EXTERN_C_INIT * @{ */ +/** + * @brief ROCProfiler External Correlation ID Operations. + * + * These kinds correspond to callback and buffered tracing kinds (@see + * ::rocprofiler_callback_tracing_kind_t and ::rocprofiler_buffer_tracing_kind_t) which generate + * correlation IDs. Typically, rocprofiler-sdk uses the most recent external correlation ID on the + * current thread set via ::rocprofiler_push_external_correlation_id; however, this approach can be + * problematic if a new external correlation ID should be set before the + * ::ROCPROFILER_CALLBACK_PHASE_ENTER callback or if relevant external correlation IDs are desired + * when the buffered tracing is used. Thus, rocprofiler-sdk provides a way for tools to get a + * callback whenever an external correlation ID is needed. However, this can add significant + * overhead for those who only need these callbacks for, say, kernel dispatches while the HSA API is + * being traced (i.e. lots of callbacks for HSA API functions). The enumeration below is provided to + * ensure that tools can default to using the external correlation IDs set via the push/pop methods + * where the external correlation ID value is not important while also getting a request for an + * external correlation ID for other tracing kinds. + */ +typedef enum // NOLINT(performance-enum-size) +{ + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_NONE = 0, ///< Unknown kind + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_CORE_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_AMD_EXT_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_IMAGE_EXT_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_FINALIZE_EXT_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HIP_RUNTIME_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HIP_COMPILER_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MARKER_CORE_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MARKER_CONTROL_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MARKER_NAME_API, ///< + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_SCRATCH_MEMORY, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_LAST, +} rocprofiler_external_correlation_id_request_kind_t; + +/** + * @brief Callback requesting a value for the external correlation id. + * + * @param [in] thread_id Id of the thread making the request + * @param [in] context_id Id of the context making the request + * @param [in] kind Origin of the external correlation id request + * @param [in] operation Regardless of whether callback or buffer tracing is being used, the + * operation value will be the same, i.e., regardless of whether callback kind is + * ::ROCPROFILER_CALLBACK_TRACING_HSA_CORE_API or the buffer record kind is + * ::ROCPROFILER_BUFFER_TRACING_HSA_CORE_API, the data/record for `hsa_init` will have an operation + * value of ::ROCPROFILER_HSA_CORE_API_ID_hsa_init. + * @param [in] internal_corr_id_value Current internal correlation ID value for the request + * @param [out] external_corr_id_value Set this value to the desired external correlation ID value + * @param [in] data The `callback_args` value passed to + * ::rocprofiler_configure_external_correlation_id_request_service. + * @returns int + * @retval 0 Used to indicate the tool had zero issues setting the external correlation ID field + * @retval 1 (or any other non-zero number) Used to indicate the callback did not set an external + * correlation ID value and the thread-local value for the most recently pushed external correlation + * ID should be used instead + */ +typedef int (*rocprofiler_external_correlation_id_request_cb_t)( + rocprofiler_thread_id_t thread_id, + rocprofiler_context_id_t context_id, + rocprofiler_external_correlation_id_request_kind_t kind, + uint32_t operation, + uint64_t internal_corr_id_value, + rocprofiler_user_data_t* external_corr_id_value, + void* data); + +/** + * @brief Configure External Correlation ID Request Service. + * + * @param [in] context_id Context to associate the service with + * @param [in] kinds Array of ::rocprofiler_external_correlation_id_request_kind_t values. If + * this parameter is null, all tracing operations will invoke the callback to request an external + * correlation ID. + * @param [in] kinds_count If the kinds array is non-null, set this to the size of the + * array. + * @param [in] callback The function to invoke for an external correlation ID request + * @param [in] callback_args Data provided to every invocation of the callback function + * @return ::rocprofiler_status_t + * @retval ::ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED Invoked outside of the initialization + * function in @ref rocprofiler_tool_configure_result_t provided to rocprofiler via @ref + * rocprofiler_configure function + * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND The provided context is not valid/registered + * @retval ::ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED if the same @ref + * rocprofiler_callback_tracing_kind_t value is provided more than once (per context) -- in + * other words, we do not support overriding or combining the kinds in separate function calls. + */ +rocprofiler_status_t +rocprofiler_configure_external_correlation_id_request_service( + rocprofiler_context_id_t context_id, + rocprofiler_external_correlation_id_request_kind_t* kinds, + size_t kinds_count, + rocprofiler_external_correlation_id_request_cb_t callback, + void* callback_args) ROCPROFILER_API ROCPROFILER_NONNULL(4); + /** * @brief Push default value for `external` field in @ref rocprofiler_correlation_id_t onto stack. * @@ -52,10 +145,11 @@ ROCPROFILER_EXTERN_C_INIT * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND Context does not exist * @retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT Thread id is not valid */ -rocprofiler_status_t ROCPROFILER_API +rocprofiler_status_t rocprofiler_push_external_correlation_id(rocprofiler_context_id_t context, rocprofiler_thread_id_t tid, - rocprofiler_user_data_t external_correlation_id); + rocprofiler_user_data_t external_correlation_id) + ROCPROFILER_API; /** * @brief Pop default value for `external` field in @ref rocprofiler_correlation_id_t off of stack @@ -67,10 +161,11 @@ rocprofiler_push_external_correlation_id(rocprofiler_context_id_t context, * @retval ::ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND Context does not exist * @retval ::ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT Thread id is not valid */ -rocprofiler_status_t ROCPROFILER_API +rocprofiler_status_t rocprofiler_pop_external_correlation_id(rocprofiler_context_id_t context, rocprofiler_thread_id_t tid, - rocprofiler_user_data_t* external_correlation_id); + rocprofiler_user_data_t* external_correlation_id) + ROCPROFILER_API; /** @} */ diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h index f61ab37548..bafd8deb15 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/fwd.h @@ -202,6 +202,52 @@ typedef enum // NOLINT(performance-enum-size) ROCPROFILER_MEMORY_COPY_LAST, } rocprofiler_memory_copy_operation_t; +/** + * @brief ROCProfiler Kernel Dispatch Tracing Operation Types. + */ +typedef enum // NOLINT(performance-enum-size) +{ + ROCPROFILER_KERNEL_DISPATCH_NONE = 0, ///< Unknown kernel dispatch operation + ROCPROFILER_KERNEL_DISPATCH_ENQUEUE = 1, + ROCPROFILER_KERNEL_DISPATCH_COMPLETE, + ROCPROFILER_KERNEL_DISPATCH_LAST, + + /// @var ROCPROFILER_KERNEL_DISPATCH_ENQUEUE + /// @brief Invoke callback prior to a kernel being enqueued and after the kernel has been + /// enqueued. When the phase is ::ROCPROFILER_CALLBACK_PHASE_ENTER, this is an opportunity to + /// push an external correlation id and/or modify the active contexts before a kernel is + /// launched. Any active contexts containing services related to a kernel dispatch (kernel + /// tracing, counter collection, etc.) will be captured after this callback and attached to the + /// kernel. These captured contexts will be considered "active" when the kernel completes even + /// if the context was stopped before the kernel completes -- this contract is designed to + /// ensure that tools do not have to delay stopping a context because of an async operation in + /// order to get the data they requested when the async operation was started. When the phase is + /// ::ROCPROFILER_CALLBACK_PHASE_EXIT, the active contexts for the kernel dispatch have been + /// captured and it is safe to disable those contexts without affecting the delivery of the + /// requested data when the kernel completes. It is important to note that, even if the context + /// associated with the kernel dispatch callback tracing service is disabled in between the + /// enter and exit phase, the exit phase callback is still delievered but that context will not + /// be captured when the kernel is enqueued and therefore will not provide a + /// ::ROCPROFILER_KERNEL_DISPATCH_COMPLETE callback. Furthermore, it should be + /// noted that if a tool encodes information into the `::rocprofiler_user_data_t` output + /// parameter in ::rocprofiler_callback_tracing_cb_t, that same value will be delivered in the + /// exit phase and in the ::ROCPROFILER_KERNEL_DISPATCH_COMPLETE callback. In + /// other words, any modifications to that user data value in the exit phase will not be + /// reflected in the ::ROCPROFILER_KERNEL_DISPATCH_COMPLETE callback because a + /// copy of that user data struct is attached to the kernel, not a reference to the user data + /// struct. + /// + /// @var ROCPROFILER_KERNEL_DISPATCH_COMPLETE + /// @brief Invoke callback after a kernel has completed and the HSA runtime has processed the + /// signal indicating that the kernel has completed. The latter half of this statement is + /// important. There is no guarantee that these callbacks are invoked in any order related to + /// when the kernels were dispatched, i.e. even if kernel A is launched and fully executed + /// before kernel B is launched, it is entirely possible that the HSA runtime ends up processing + /// the signal associated with kernel B before processing the signal associated with kernel A -- + /// resulting in rocprofiler-sdk invoking this operation callback for kernel B before invoking + /// the callback for kernel A. +} rocprofiler_kernel_dispatch_operation_t; + /** * @brief PC Sampling Method. */ @@ -335,6 +381,12 @@ typedef uint32_t rocprofiler_tracing_operation_t; */ typedef uint64_t rocprofiler_kernel_id_t; +// /** +// * @brief Sequence identifier type +// * +// */ +typedef uint64_t rocprofiler_dispatch_id_t; + // forward declaration of struct typedef struct rocprofiler_pc_sampling_configuration_s rocprofiler_pc_sampling_configuration_t; typedef struct rocprofiler_pc_sampling_record_s rocprofiler_pc_sampling_record_t; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/CMakeLists.txt index 59f50b498a..d2e5ea7ab8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/CMakeLists.txt @@ -43,6 +43,8 @@ add_subdirectory(aql) add_subdirectory(pc_sampling) add_subdirectory(marker) add_subdirectory(thread_trace) +add_subdirectory(tracing) +add_subdirectory(kernel_dispatch) target_link_libraries( rocprofiler-object-library diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/agent.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/agent.cpp index 1d666bf17a..f00b8a51ce 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/agent.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/agent.cpp @@ -52,22 +52,14 @@ #include #include -namespace fs = rocprofiler::common::filesystem; - -#if defined(ROCPROFILER_CI) -# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__) -# define ROCP_CI_LOG(NON_CI_LEVEL, ...) ROCP_FATAL -#else -# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__) -# define ROCP_CI_LOG(NON_CI_LEVEL, ...) LOG(NON_CI_LEVEL) -#endif - namespace rocprofiler { namespace agent { namespace { +namespace fs = ::rocprofiler::common::filesystem; + using name_array_t = std::vector>>; name_array_t* @@ -813,7 +805,7 @@ construct_agent_cache(::HsaApiTable* table) } } - ROCP_ERROR << "# agent node maps: " << hsa_agent_node_map.size(); + ROCP_INFO << "# agent node maps: " << hsa_agent_node_map.size(); LOG_IF(FATAL, agent_map.size() != hsa_agents.size()) << "rocprofiler was only able to map " << agent_map.size() @@ -905,10 +897,8 @@ construct_agent_cache(::HsaApiTable* table) std::optional get_hsa_agent(const rocprofiler_agent_t* agent) { - ROCP_ERROR << "# of agent mappings: " << get_agent_mapping().size(); for(const auto& itr : get_agent_mapping()) { - ROCP_ERROR << "checking " << itr.rocp_agent->id.handle << " vs. " << agent->id.handle; if(itr.rocp_agent->id.handle == agent->id.handle) return itr.hsa_agent; } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp index 74086d04d5..fad77dbf3f 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/buffer_tracing.cpp @@ -32,6 +32,7 @@ #include "lib/rocprofiler-sdk/hsa/async_copy.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/scratch_memory.hpp" +#include "lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp" #include "lib/rocprofiler-sdk/marker/marker.hpp" #include "lib/rocprofiler-sdk/registration.hpp" @@ -227,6 +228,10 @@ rocprofiler_query_buffer_tracing_kind_operation_name(rocprofiler_buffer_tracing_ break; } case ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH: + { + val = rocprofiler::kernel_dispatch::name_by_id(operation); + break; + } case ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION: case ROCPROFILER_BUFFER_TRACING_CORRELATION_ID_RETIREMENT: { @@ -330,6 +335,10 @@ rocprofiler_iterate_buffer_tracing_kind_operations( break; } case ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH: + { + ops = rocprofiler::kernel_dispatch::get_ids(); + break; + } case ROCPROFILER_BUFFER_TRACING_PAGE_MIGRATION: case ROCPROFILER_BUFFER_TRACING_CORRELATION_ID_RETIREMENT: { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp index d0f299569b..1efb6bb792 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/callback_tracing.cpp @@ -20,6 +20,7 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE // SOFTWARE. +#include #include #include #include @@ -32,6 +33,7 @@ #include "lib/rocprofiler-sdk/hsa/code_object.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/scratch_memory.hpp" +#include "lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp" #include "lib/rocprofiler-sdk/marker/marker.hpp" #include "lib/rocprofiler-sdk/registration.hpp" @@ -103,10 +105,6 @@ rocprofiler_configure_callback_tracing_service(rocprofiler_context_id_t if(rocprofiler::registration::get_init_status() > -1) return ROCPROFILER_STATUS_ERROR_CONFIGURATION_LOCKED; - static auto unsupported = std::unordered_set{ - ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH}; - if(unsupported.count(kind) > 0) return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; - auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id); if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; @@ -219,7 +217,7 @@ rocprofiler_query_callback_tracing_kind_operation_name(rocprofiler_callback_trac } case ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH: { - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + val = rocprofiler::kernel_dispatch::name_by_id(operation); } }; @@ -322,7 +320,8 @@ rocprofiler_iterate_callback_tracing_kind_operations( } case ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH: { - return ROCPROFILER_STATUS_ERROR_NOT_IMPLEMENTED; + ops = rocprofiler::kernel_dispatch::get_ids(); + break; } }; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp index 97de3c75c2..081e3b516a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/context/context.cpp @@ -131,7 +131,11 @@ get_unique_internal_id() uint32_t correlation_id::add_ref_count() { - return m_ref_count.fetch_add(1); + auto _ret = m_ref_count.fetch_add(1); + + LOG_IF(FATAL, _ret == 0) << "correlation id already retired"; + + return _ret; } uint32_t @@ -190,6 +194,8 @@ correlation_id::sub_kern_count() correlation_id* correlation_tracing_service::construct(uint32_t _init_ref_count) { + LOG_IF(FATAL, _init_ref_count == 0) << "must have reference count > 0"; + auto _internal_id = get_unique_internal_id(); auto* corr_id_map = get_correlation_id_map(); if(!corr_id_map) return nullptr; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp index b41ebe807d..262c8838ea 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp @@ -266,7 +266,8 @@ queue_cb(const context::context* ctx, const std::shared_ptr& info, const hsa::Queue& queue, const hsa::rocprofiler_packet& pkt, - uint64_t kernel_id, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, rocprofiler_user_data_t* user_data, const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, const context::correlation_id* correlation_id) @@ -326,6 +327,7 @@ queue_cb(const context::context* ctx, common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); dispatch_data.kernel_id = kernel_id; + dispatch_data.dispatch_id = dispatch_id; dispatch_data.agent_id = CHECK_NOTNULL(queue.get_agent().get_rocp_agent())->id; dispatch_data.queue_id = queue.get_id(); dispatch_data.correlation_id = _corr_id_v; @@ -375,7 +377,7 @@ queue_cb(const context::context* ctx, void completed_cb(const context::context* ctx, const std::shared_ptr& info, - const hsa::Queue& queue, + const hsa::Queue& /*queue*/, hsa::rocprofiler_packet, const hsa::Queue::queue_info_session_t& session, inst_pkt_t& pkts) @@ -432,8 +434,8 @@ completed_cb(const context::context* ctx, if(const auto* _corr_id = session.correlation_id) { _corr_id_v.internal = _corr_id->internal; - if(const auto* external = - rocprofiler::common::get_val(session.extern_corr_ids, info->internal_context)) + if(const auto* external = rocprofiler::common::get_val( + session.tracing_data.external_correlation_ids, info->internal_context)) { _corr_id_v.external = *external; } @@ -463,20 +465,15 @@ completed_cb(const context::context* ctx, auto dispatch_data = common::init_public_api_struct(rocprofiler_profile_counting_dispatch_data_t{}); - const auto& kernel_dispatch_pkt = session.kernel_pkt.kernel_dispatch; - - dispatch_data.kernel_id = session.kernel_id; - dispatch_data.agent_id = CHECK_NOTNULL(queue.get_agent().get_rocp_agent())->id; - dispatch_data.queue_id = queue.get_id(); + dispatch_data.kernel_id = session.callback_record.kernel_id; + dispatch_data.dispatch_id = session.callback_record.dispatch_id; + dispatch_data.agent_id = session.callback_record.agent_id; + dispatch_data.queue_id = session.callback_record.queue_id; dispatch_data.correlation_id = _corr_id_v; - dispatch_data.private_segment_size = kernel_dispatch_pkt.private_segment_size; - dispatch_data.group_segment_size = kernel_dispatch_pkt.group_segment_size; - dispatch_data.workgroup_size = {kernel_dispatch_pkt.workgroup_size_x, - kernel_dispatch_pkt.workgroup_size_y, - kernel_dispatch_pkt.workgroup_size_z}; - dispatch_data.grid_size = {kernel_dispatch_pkt.grid_size_x, - kernel_dispatch_pkt.grid_size_y, - kernel_dispatch_pkt.grid_size_z}; + dispatch_data.private_segment_size = session.callback_record.private_segment_size; + dispatch_data.group_segment_size = session.callback_record.group_segment_size; + dispatch_data.workgroup_size = session.callback_record.workgroup_size; + dispatch_data.grid_size = session.callback_record.grid_size; info->record_callback( dispatch_data, out.data(), out.size(), session.user_data, info->record_callback_args); @@ -509,7 +506,8 @@ start_context(const context::context* ctx) std::nullopt, [=](const hsa::Queue& q, const hsa::rocprofiler_packet& kern_pkt, - uint64_t kernel_id, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, rocprofiler_user_data_t* user_data, const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, const context::correlation_id* correlation_id) { @@ -518,6 +516,7 @@ start_context(const context::context* ctx) q, kern_pkt, kernel_id, + dispatch_id, user_data, extern_corr_ids, correlation_id); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp index c979aa3434..e5b07f9f43 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp @@ -130,7 +130,8 @@ queue_cb(const context::context* ctx, const std::shared_ptr& info, const hsa::Queue& queue, const hsa::rocprofiler_packet& pkt, - uint64_t kernel_id, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, rocprofiler_user_data_t* user_data, const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, const context::correlation_id* correlation_id); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp index 45f98d678d..ef6bd1721d 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/tests/core.cpp @@ -365,6 +365,7 @@ struct expected_dispatch rocprofiler_queue_id_t queue_id = {.handle = 0}; rocprofiler_agent_id_t agent_id = {.handle = 0}; uint64_t kernel_id = 0; + uint64_t dispatch_id = 0; rocprofiler_correlation_id_t correlation_id = {.internal = 0, .external = {.value = 0}}; rocprofiler_dim3_t workgroup_size = {0, 0, 0}; rocprofiler_dim3_t grid_size = {0, 0, 0}; @@ -383,9 +384,11 @@ user_dispatch_cb(rocprofiler_profile_counting_dispatch_data_t dispatch_data, auto queue_id = dispatch_data.queue_id; auto correlation_id = dispatch_data.correlation_id; auto kernel_id = dispatch_data.kernel_id; + auto dispatch_id = dispatch_data.dispatch_id; EXPECT_EQ(sizeof(rocprofiler_profile_counting_dispatch_data_t), dispatch_data.size); EXPECT_EQ(expected.kernel_id, kernel_id); + EXPECT_EQ(expected.dispatch_id, dispatch_id); EXPECT_EQ(expected.agent_id, agent_id); EXPECT_EQ(expected.queue_id.handle, queue_id.handle); EXPECT_EQ(expected.correlation_id.internal, correlation_id.internal); @@ -480,14 +483,22 @@ TEST(core, check_callbacks) pkt.kernel_dispatch.grid_size_y, pkt.kernel_dispatch.grid_size_z}; expected.kernel_id = count++; + expected.dispatch_id = count++; expected.queue_id = qid; expected.agent_id = fq.get_agent().get_rocp_agent()->id; hsa::Queue::queue_info_session_t::external_corr_id_map_t extern_ids = {}; auto user_data = rocprofiler_user_data_t{.value = corr_id.internal}; - auto ret_pkt = counters::queue_cb( - &ctx, cb_info, fq, pkt, expected.kernel_id, &user_data, extern_ids, &corr_id); + auto ret_pkt = counters::queue_cb(&ctx, + cb_info, + fq, + pkt, + expected.kernel_id, + expected.dispatch_id, + &user_data, + extern_ids, + &corr_id); ASSERT_TRUE(ret_pkt) << fmt::format("Expected a packet to be generated for - {}", metric.name()); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.cpp index c2b3601e61..8646f9f55a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.cpp @@ -77,6 +77,32 @@ external_correlation::get(rocprofiler_thread_id_t tid) const tid); } +rocprofiler_user_data_t +external_correlation::get(rocprofiler_thread_id_t tid, + const context::context* ctx, + request_kind_t kind, + uint32_t op, + uint64_t internal_corr_id) const +{ + if(requires_request(kind)) + { + auto opt_data = invoke_callback(tid, ctx, kind, op, internal_corr_id); + if(opt_data) return *opt_data; + } + + return get(tid); +} + +rocprofiler_user_data_t& +external_correlation::update(rocprofiler_user_data_t& value, + rocprofiler_thread_id_t thr_id, + request_kind_t kind) const +{ + // if requires request is true, do not update, otherwise, get the latest pushed external + // correlation id + return (requires_request(kind)) ? value : (value = get(thr_id)); +} + void external_correlation::push(rocprofiler_thread_id_t tid, rocprofiler_user_data_t user_data) { @@ -139,10 +165,90 @@ external_correlation::pop(rocprofiler_thread_id_t tid) }, tid); } + +rocprofiler_status_t +external_correlation::configure_request(request_cb_t callback_v, + void* callback_data_v, + const std::vector& kinds_v) +{ + if(!callback_v) + return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + else if(callback || callback_data || request.any()) + return ROCPROFILER_STATUS_ERROR_SERVICE_ALREADY_CONFIGURED; + + callback = callback_v; + callback_data = callback_data_v; + + if(kinds_v.empty()) + { + request.flip(); + } + else + { + for(auto itr : kinds_v) + { + if(itr <= ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_NONE || + itr >= ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_LAST) + return ROCPROFILER_STATUS_ERROR_KIND_NOT_FOUND; + + request.set(itr - 1, true); + } + } + + return ROCPROFILER_STATUS_SUCCESS; +} + +bool +external_correlation::requires_request(request_kind_t kind) const +{ + return request.test(kind - 1); +} + +std::optional +external_correlation::invoke_callback(rocprofiler_thread_id_t thr_id, + const context::context* ctx, + request_kind_t kind, + uint32_t op, + uint64_t internal_corr_id) const +{ + auto value = rocprofiler_user_data_t{.value = 0}; + auto ctx_id = rocprofiler_context_id_t{ctx->context_idx}; + + if(callback(thr_id, ctx_id, kind, op, internal_corr_id, &value, callback_data) == 0) + return value; + + return std::nullopt; +} } // namespace external_correlation } // namespace rocprofiler extern "C" { +rocprofiler_status_t +rocprofiler_configure_external_correlation_id_request_service( + rocprofiler_context_id_t context_id, + rocprofiler_external_correlation_id_request_kind_t* kinds, + size_t kinds_count, + rocprofiler_external_correlation_id_request_cb_t callback, + void* callback_args) +{ + auto* ctx = rocprofiler::context::get_mutable_registered_context(context_id); + if(!ctx) return ROCPROFILER_STATUS_ERROR_CONTEXT_NOT_FOUND; + + auto kinds_v = std::vector{}; + if(kinds) + { + kinds_v.reserve(kinds_count); + + if(kinds_count < 1) return ROCPROFILER_STATUS_ERROR_INVALID_ARGUMENT; + + for(size_t i = 0; i < kinds_count; ++i) + kinds_v.emplace_back(kinds[i]); + } + + return ctx->correlation_tracer.external_correlator.configure_request( + callback, callback_args, kinds_v); +} + rocprofiler_status_t rocprofiler_push_external_correlation_id(rocprofiler_context_id_t context, rocprofiler_thread_id_t tid, diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.hpp index 981d96895a..3bde84c3a4 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/external_correlation.hpp @@ -29,13 +29,19 @@ #include "lib/common/synchronized.hpp" #include "lib/common/utility.hpp" +#include #include +#include #include #include #include namespace rocprofiler { +namespace context +{ +struct context; +} namespace external_correlation { static constexpr bool enable_const_wlock_v = true; @@ -50,12 +56,43 @@ using external_correlation_map_t = struct external_correlation { - rocprofiler_user_data_t get(rocprofiler_thread_id_t) const; + using request_cb_t = rocprofiler_external_correlation_id_request_cb_t; + using request_kind_t = rocprofiler_external_correlation_id_request_kind_t; + + static constexpr size_t request_kind_size = ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_LAST - 1; + + rocprofiler_user_data_t get(rocprofiler_thread_id_t thr_id, + const context::context* ctx, + request_kind_t kind, + uint32_t op, + uint64_t internal_corr_id) const; + rocprofiler_user_data_t& update(rocprofiler_user_data_t& data, + rocprofiler_thread_id_t thr_id, + request_kind_t kind) const; + void push(rocprofiler_thread_id_t, rocprofiler_user_data_t); rocprofiler_user_data_t pop(rocprofiler_thread_id_t); + rocprofiler_status_t configure_request(request_cb_t callback_v, + void* callback_data_v, + const std::vector& kinds_v); + + bool requires_request(request_kind_t kind) const; + private: - common::Synchronized data = {}; + rocprofiler_user_data_t get(rocprofiler_thread_id_t thr_id) const; + + std::optional invoke_callback( + rocprofiler_thread_id_t thr_id, + const context::context* ctx, + rocprofiler_external_correlation_id_request_kind_t kind, + uint32_t op, + uint64_t internal_corr_id) const; + + request_cb_t callback = nullptr; + void* callback_data = nullptr; + std::bitset request = 0; + common::Synchronized data = {}; }; } // namespace external_correlation } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp index 64fef33c6c..7d5dae3069 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.cpp @@ -29,6 +29,7 @@ #include "lib/rocprofiler-sdk/hip/types.hpp" #include "lib/rocprofiler-sdk/hip/utils.hpp" #include "lib/rocprofiler-sdk/registration.hpp" +#include "lib/rocprofiler-sdk/tracing/tracing.hpp" #include #include @@ -179,57 +180,6 @@ hip_api_impl::exec(FuncT&& _func, Args&&... args) return get_default_retval(); } -namespace -{ -using correlation_service = context::correlation_tracing_service; - -struct callback_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_callback_tracing_record_t record = {}; - rocprofiler_user_data_t user_data = {.value = 0}; -}; - -struct buffered_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_user_data_t external_correlation = {}; -}; - -constexpr auto empty_user_data = rocprofiler_user_data_t{.value = 0}; - -void -populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, - rocprofiler_buffer_tracing_kind_t buffered_domain_idx, - int operation_idx, - std::vector& callback_contexts, - std::vector& buffered_contexts) -{ - auto active_contexts = context::context_array_t{}; - auto thr_id = common::get_tid(); - for(const auto* itr : context::get_active_contexts(active_contexts)) - { - if(!itr) continue; - - if(itr->callback_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->callback_tracer->domains(callback_domain_idx, operation_idx)) - callback_contexts.emplace_back( - callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); - } - - if(itr->buffered_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->buffered_tracer->domains(buffered_domain_idx, operation_idx)) - buffered_contexts.emplace_back(buffered_context_data{ - itr, itr->correlation_tracer.external_correlator.get(thr_id)}); - } - } -} -} // namespace - template template RetT @@ -239,17 +189,23 @@ hip_api_impl::functor(Args... args) using callback_api_data_t = typename hip_domain_info::callback_data_type; using buffered_api_data_t = typename hip_domain_info::buffered_data_type; + constexpr auto external_corr_id_domain_idx = + hip_domain_info::external_correlation_id_domain_idx; + LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name; - auto thr_id = common::get_tid(); - auto callback_contexts = std::vector{}; - auto buffered_contexts = std::vector{}; + constexpr auto ref_count = 2; + auto thr_id = common::get_tid(); + auto callback_contexts = tracing::callback_context_data_vec_t{}; + auto buffered_contexts = tracing::buffered_context_data_vec_t{}; + auto external_corr_ids = tracing::external_correlation_id_map_t{}; - populate_contexts(info_type::callback_domain_idx, - info_type::buffered_domain_idx, - info_type::operation_idx, - callback_contexts, - buffered_contexts); + tracing::populate_contexts(info_type::callback_domain_idx, + info_type::buffered_domain_idx, + info_type::operation_idx, + callback_contexts, + buffered_contexts, + external_corr_ids); if(callback_contexts.empty() && buffered_contexts.empty()) { @@ -260,25 +216,19 @@ hip_api_impl::functor(Args... args) return; } - constexpr auto ref_count = 2; - auto buffer_record = common::init_public_api_struct(buffered_api_data_t{}); - auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); - auto* corr_id = correlation_service::construct(ref_count); - auto internal_corr_id = corr_id->internal; + LOG_IF(FATAL, external_corr_ids.size() < (callback_contexts.size() + buffered_contexts.size())) + << "missing external correlation ids"; - // construct the buffered info before the callback so the callbacks are as closely wrapped - // around the function call as possible - if(!buffered_contexts.empty()) - { - buffer_record.kind = info_type::buffered_domain_idx; - // external correlation will be updated right before record is placed in buffer - buffer_record.correlation_id = - rocprofiler_correlation_id_t{internal_corr_id, empty_user_data}; - buffer_record.operation = info_type::operation_idx; - buffer_record.thread_id = thr_id; - } + auto buffer_record = common::init_public_api_struct(buffered_api_data_t{}); + auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); + auto* corr_id = tracing::correlation_service::construct(ref_count); + auto internal_corr_id = corr_id->internal; - tracer_data.size = sizeof(callback_api_data_t); + tracing::populate_external_correlation_ids(external_corr_ids, + thr_id, + external_corr_id_domain_idx, + info_type::operation_idx, + internal_corr_id); // invoke the callbacks if(!callback_contexts.empty()) @@ -286,42 +236,22 @@ hip_api_impl::functor(Args... args) set_data_args(info_type::get_api_data_args(tracer_data.args), convert_arg_type(std::forward(args))...); - for(auto& itr : callback_contexts) - { - auto& ctx = itr.ctx; - auto& record = itr.record; - auto& user_data = itr.user_data; - - auto extern_corr_id_v = ctx->correlation_tracer.external_correlator.get(thr_id); - - auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; - record = - rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, - thr_id, - corr_id_v, - info_type::callback_domain_idx, - info_type::operation_idx, - ROCPROFILER_CALLBACK_PHASE_ENTER, - static_cast(&tracer_data)}; - - auto& callback_info = - ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); - callback_info.callback(record, &user_data, callback_info.data); - - // enter callback may update the external correlation id field - record.correlation_id.external = - ctx->correlation_tracer.external_correlator.get(thr_id); - } + tracing::execute_phase_enter_callbacks(callback_contexts, + thr_id, + internal_corr_id, + external_corr_ids, + info_type::callback_domain_idx, + info_type::operation_idx, + tracer_data); } + // enter callback may update the external correlation id field + tracing::update_external_correlation_ids( + external_corr_ids, thr_id, external_corr_id_domain_idx); + // record the start timestamp as close to the function call as possible if(!buffered_contexts.empty()) { - for(auto& itr : buffered_contexts) - { - itr.external_correlation = itr.ctx->correlation_tracer.external_correlator.get(thr_id); - } - buffer_record.start_timestamp = common::timestamp_ns(); } @@ -340,41 +270,22 @@ hip_api_impl::functor(Args... args) { set_data_retval(tracer_data.retval, _ret); - for(auto& itr : callback_contexts) - { - auto& ctx = itr.ctx; - auto& record = itr.record; - auto& user_data = itr.user_data; - - record.phase = ROCPROFILER_CALLBACK_PHASE_EXIT; - record.payload = static_cast(&tracer_data); - - auto& callback_info = - ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); - callback_info.callback(record, &user_data, callback_info.data); - } + tracing::execute_phase_exit_callbacks(callback_contexts, + external_corr_ids, + info_type::callback_domain_idx, + info_type::operation_idx, + tracer_data); } if(!buffered_contexts.empty()) { - for(auto& itr : buffered_contexts) - { - assert(itr.ctx->buffered_tracer); - auto buffer_id = - itr.ctx->buffered_tracer->buffer_data.at(info_type::buffered_domain_idx); - auto buffer_v = buffer::get_buffer(buffer_id); - if(buffer_v && buffer_v->context_id == itr.ctx->context_idx && - buffer_v->buffer_id == buffer_id.handle) - { - // make copy of record - auto record_v = buffer_record; - // update the record with the correlation - record_v.correlation_id.external = itr.external_correlation; - - buffer_v->emplace( - ROCPROFILER_BUFFER_CATEGORY_TRACING, info_type::buffered_domain_idx, record_v); - } - } + tracing::execute_buffer_record_emplace(buffered_contexts, + thr_id, + internal_corr_id, + external_corr_ids, + info_type::buffered_domain_idx, + info_type::operation_idx, + buffer_record); } // decrement the reference count after usage in the callback/buffers diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp index 2232ca110b..9e7d485e47 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hip/hip.def.cpp @@ -23,6 +23,7 @@ #if defined(ROCPROFILER_LIB_ROCPROFILER_HIP_HIP_CPP_IMPL) && \ ROCPROFILER_LIB_ROCPROFILER_HIP_HIP_CPP_IMPL == 1 +# include # include # include "defines.hpp" @@ -49,6 +50,8 @@ struct hip_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_HIP_RUNTIME_API; static constexpr auto none = ROCPROFILER_HIP_RUNTIME_API_ID_NONE; static constexpr auto last = ROCPROFILER_HIP_RUNTIME_API_ID_LAST; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HIP_RUNTIME_API; }; template <> @@ -59,6 +62,8 @@ struct hip_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_HIP_COMPILER_API; static constexpr auto none = ROCPROFILER_HIP_COMPILER_API_ID_NONE; static constexpr auto last = ROCPROFILER_HIP_COMPILER_API_ID_LAST; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HIP_COMPILER_API; }; } // namespace hip } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt index c9e0927fd6..74a19cd342 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/CMakeLists.txt @@ -4,11 +4,11 @@ set(ROCPROFILER_LIB_HSA_SOURCES aql_packet.cpp async_copy.cpp code_object.cpp + hsa_barrier.cpp hsa.cpp + profile_serializer.cpp queue_controller.cpp queue.cpp - hsa_barrier.cpp - profile_serializer.cpp scratch_memory.cpp) set(ROCPROFILER_LIB_HSA_HEADERS @@ -22,6 +22,8 @@ set(ROCPROFILER_LIB_HSA_HEADERS profile_serializer.hpp queue_controller.hpp queue.hpp + queue_info_session.hpp + rocprofiler_packet.hpp scratch_memory.hpp types.hpp utils.hpp) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp index bf8c70df40..735390a72a 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp @@ -22,6 +22,7 @@ #include "lib/rocprofiler-sdk/hsa/async_copy.hpp" #include "lib/common/defines.hpp" +#include "lib/common/scope_destructor.hpp" #include "lib/common/static_object.hpp" #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/agent.hpp" @@ -70,21 +71,6 @@ namespace rocprofiler { namespace hsa { -namespace -{ -std::string_view -get_hsa_status_string(hsa_status_t _status) -{ - if(!hsa::get_core_table()->hsa_status_string_fn) return std::string_view{"(unknown HSA error)"}; - const char* _status_msg = nullptr; - return (hsa::get_core_table()->hsa_status_string_fn(_status, &_status_msg) == - HSA_STATUS_SUCCESS && - _status_msg) - ? std::string_view{_status_msg} - : std::string_view{"(unknown HSA error)"}; -} -} // namespace - namespace async_copy { namespace @@ -598,18 +584,39 @@ async_copy_impl(Args... args) } } - _data->correlation_id = context::get_latest_correlation_id(); - auto& extern_corr_ids = _data->extern_corr_ids; + _data->correlation_id = context::get_latest_correlation_id(); + context::correlation_id* _corr_id_pop = nullptr; + + if(!_data->correlation_id) + { + constexpr auto ref_count = 1; + _data->correlation_id = context::correlation_tracing_service::construct(ref_count); + _corr_id_pop = _data->correlation_id; + } // increase the reference count to denote that this correlation id is being used in a kernel - if(_data->correlation_id) - { - extern_corr_ids.reserve(_data->contexts.size()); // reserve for performance - for(const auto* ctx : _data->contexts) - extern_corr_ids.emplace(ctx, - ctx->correlation_tracer.external_correlator.get(_data->tid)); - _data->correlation_id->add_ref_count(); - } + _data->correlation_id->add_ref_count(); + + // if we constructed a correlation id, this decrements the reference count after the underlying + // function returns + auto _corr_id_dtor = common::scope_destructor{[_corr_id_pop]() { + if(_corr_id_pop) + { + context::pop_latest_correlation_id(_corr_id_pop); + _corr_id_pop->sub_ref_count(); + } + }}; + + auto& extern_corr_ids = _data->extern_corr_ids; + extern_corr_ids.reserve(_data->contexts.size()); // reserve for performance + for(const auto* ctx : _data->contexts) + extern_corr_ids.emplace(ctx, + ctx->correlation_tracer.external_correlator.get( + _data->tid, + ctx, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MEMORY_COPY, + _direction, + _data->correlation_id->internal)); _data->orig_signal = _completion_signal; _completion_signal = _data->rocp_signal; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp index 46ba89bdda..c31249f245 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.cpp @@ -31,6 +31,7 @@ #include "lib/rocprofiler-sdk/hsa/types.hpp" #include "lib/rocprofiler-sdk/hsa/utils.hpp" #include "lib/rocprofiler-sdk/registration.hpp" +#include "lib/rocprofiler-sdk/tracing/tracing.hpp" #include #include @@ -258,65 +259,17 @@ hsa_api_impl::exec(FuncT&& _func, Args&&... args) return return_type{HSA_STATUS_ERROR}; } -namespace -{ -using correlation_service = context::correlation_tracing_service; -using buffer_hsa_api_record_t = rocprofiler_buffer_tracing_hsa_api_record_t; -using callback_hsa_api_data_t = rocprofiler_callback_tracing_hsa_api_data_t; - -struct callback_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_callback_tracing_record_t record = {}; - rocprofiler_user_data_t user_data = {.value = 0}; -}; - -struct buffered_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_user_data_t external_correlation = {}; -}; - -constexpr auto empty_user_data = rocprofiler_user_data_t{.value = 0}; - -void -populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, - rocprofiler_buffer_tracing_kind_t buffered_domain_idx, - int operation_idx, - std::vector& callback_contexts, - std::vector& buffered_contexts) -{ - auto active_contexts = context::context_array_t{}; - auto thr_id = common::get_tid(); - for(const auto* itr : context::get_active_contexts(active_contexts)) - { - if(!itr) continue; - - if(itr->callback_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->callback_tracer->domains(callback_domain_idx, operation_idx)) - callback_contexts.emplace_back( - callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); - } - - if(itr->buffered_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->buffered_tracer->domains(buffered_domain_idx, operation_idx)) - buffered_contexts.emplace_back(buffered_context_data{ - itr, itr->correlation_tracer.external_correlator.get(thr_id)}); - } - } -} -} // namespace - template template RetT hsa_api_impl::functor(Args... args) { - using info_type = hsa_api_info; + using buffer_hsa_api_record_t = rocprofiler_buffer_tracing_hsa_api_record_t; + using callback_hsa_api_data_t = rocprofiler_callback_tracing_hsa_api_data_t; + using info_type = hsa_api_info; + + constexpr auto external_corr_id_domain_idx = + hsa_domain_info::external_correlation_id_domain_idx; LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name; @@ -329,15 +282,18 @@ hsa_api_impl::functor(Args... args) return; } - auto thr_id = common::get_tid(); - auto callback_contexts = std::vector{}; - auto buffered_contexts = std::vector{}; + constexpr auto ref_count = 2; + auto thr_id = common::get_tid(); + auto callback_contexts = tracing::callback_context_data_vec_t{}; + auto buffered_contexts = tracing::buffered_context_data_vec_t{}; + auto external_corr_ids = tracing::external_correlation_id_map_t{}; - populate_contexts(info_type::callback_domain_idx, - info_type::buffered_domain_idx, - info_type::operation_idx, - callback_contexts, - buffered_contexts); + tracing::populate_contexts(info_type::callback_domain_idx, + info_type::buffered_domain_idx, + info_type::operation_idx, + callback_contexts, + buffered_contexts, + external_corr_ids); if(callback_contexts.empty() && buffered_contexts.empty()) { @@ -348,68 +304,41 @@ hsa_api_impl::functor(Args... args) return; } - constexpr auto ref_count = 2; - auto buffer_record = common::init_public_api_struct(buffer_hsa_api_record_t{}); - auto tracer_data = common::init_public_api_struct(callback_hsa_api_data_t{}); - auto* corr_id = correlation_service::construct(ref_count); - auto internal_corr_id = corr_id->internal; + LOG_IF(FATAL, external_corr_ids.size() < (callback_contexts.size() + buffered_contexts.size())) + << "missing external correlation ids"; - // construct the buffered info before the callback so the callbacks are as closely wrapped - // around the function call as possible - if(!buffered_contexts.empty()) - { - buffer_record.kind = info_type::buffered_domain_idx; - // external correlation will be updated right before record is placed in buffer - buffer_record.correlation_id = - rocprofiler_correlation_id_t{internal_corr_id, empty_user_data}; - buffer_record.operation = info_type::operation_idx; - buffer_record.thread_id = thr_id; - } + auto buffer_record = common::init_public_api_struct(buffer_hsa_api_record_t{}); + auto tracer_data = common::init_public_api_struct(callback_hsa_api_data_t{}); + auto* corr_id = tracing::correlation_service::construct(ref_count); + auto internal_corr_id = corr_id->internal; - tracer_data.size = sizeof(rocprofiler_callback_tracing_hsa_api_data_t); - set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward(args)...); + tracing::populate_external_correlation_ids(external_corr_ids, + thr_id, + external_corr_id_domain_idx, + info_type::operation_idx, + internal_corr_id); // invoke the callbacks if(!callback_contexts.empty()) { set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward(args)...); - for(auto& itr : callback_contexts) - { - auto& ctx = itr.ctx; - auto& record = itr.record; - auto& user_data = itr.user_data; - - auto extern_corr_id_v = ctx->correlation_tracer.external_correlator.get(thr_id); - - auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; - record = - rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, - thr_id, - corr_id_v, - info_type::callback_domain_idx, - info_type::operation_idx, - ROCPROFILER_CALLBACK_PHASE_ENTER, - static_cast(&tracer_data)}; - - auto& callback_info = - ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); - callback_info.callback(record, &user_data, callback_info.data); - - // enter callback may update the external correlation id field - record.correlation_id.external = - ctx->correlation_tracer.external_correlator.get(thr_id); - } + tracing::execute_phase_enter_callbacks(callback_contexts, + thr_id, + internal_corr_id, + external_corr_ids, + info_type::callback_domain_idx, + info_type::operation_idx, + tracer_data); } + // enter callback may update the external correlation id field + tracing::update_external_correlation_ids( + external_corr_ids, thr_id, external_corr_id_domain_idx); + // record the start timestamp as close to the function call as possible if(!buffered_contexts.empty()) { - for(auto& itr : buffered_contexts) - { - itr.external_correlation = itr.ctx->correlation_tracer.external_correlator.get(thr_id); - } - buffer_record.start_timestamp = common::timestamp_ns(); } @@ -428,41 +357,22 @@ hsa_api_impl::functor(Args... args) { set_data_retval(tracer_data.retval, _ret); - for(auto& itr : callback_contexts) - { - auto& ctx = itr.ctx; - auto& record = itr.record; - auto& user_data = itr.user_data; - - record.phase = ROCPROFILER_CALLBACK_PHASE_EXIT; - record.payload = static_cast(&tracer_data); - - auto& callback_info = - ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); - callback_info.callback(record, &user_data, callback_info.data); - } + tracing::execute_phase_exit_callbacks(callback_contexts, + external_corr_ids, + info_type::callback_domain_idx, + info_type::operation_idx, + tracer_data); } if(!buffered_contexts.empty()) { - for(auto& itr : buffered_contexts) - { - assert(itr.ctx->buffered_tracer); - auto buffer_id = - itr.ctx->buffered_tracer->buffer_data.at(info_type::buffered_domain_idx); - auto buffer_v = buffer::get_buffer(buffer_id); - if(buffer_v && buffer_v->context_id == itr.ctx->context_idx && - buffer_v->buffer_id == buffer_id.handle) - { - // make copy of record - auto record_v = buffer_record; - // update the record with the correlation - record_v.correlation_id.external = itr.external_correlation; - - buffer_v->emplace( - ROCPROFILER_BUFFER_CATEGORY_TRACING, info_type::buffered_domain_idx, record_v); - } - } + tracing::execute_buffer_record_emplace(buffered_contexts, + thr_id, + internal_corr_id, + external_corr_ids, + info_type::buffered_domain_idx, + info_type::operation_idx, + buffer_record); } // decrement the reference count after usage in the callback/buffers @@ -692,6 +602,17 @@ update_table(const context::context_array_t& _contexts, } } // namespace +std::string_view +get_hsa_status_string(hsa_status_t _status) +{ + const char* _status_msg = nullptr; + return (hsa::get_core_table()->hsa_status_string_fn(_status, &_status_msg) == + HSA_STATUS_SUCCESS && + _status_msg) + ? std::string_view{_status_msg} + : std::string_view{"(unknown HSA error)"}; +} + // check out the assembly here... this compiles to a switch statement template const char* diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp index c120b39500..a5f27ea6cd 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.def.cpp @@ -24,6 +24,7 @@ #include "lib/rocprofiler-sdk/hsa/defines.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" +#include #include #include #include @@ -58,6 +59,8 @@ struct hsa_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_HSA_CORE_API; static constexpr auto none = ROCPROFILER_HSA_CORE_API_ID_NONE; static constexpr auto last = ROCPROFILER_HSA_CORE_API_ID_LAST; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_CORE_API; }; template <> @@ -69,6 +72,8 @@ struct hsa_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_HSA_AMD_EXT_API; static constexpr auto none = ROCPROFILER_HSA_AMD_EXT_API_ID_NONE; static constexpr auto last = ROCPROFILER_HSA_AMD_EXT_API_ID_LAST; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_AMD_EXT_API; }; template <> @@ -80,6 +85,8 @@ struct hsa_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_HSA_IMAGE_EXT_API; static constexpr auto none = ROCPROFILER_HSA_IMAGE_EXT_API_ID_NONE; static constexpr auto last = ROCPROFILER_HSA_IMAGE_EXT_API_ID_LAST; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_IMAGE_EXT_API; }; template <> @@ -91,6 +98,8 @@ struct hsa_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_HSA_FINALIZE_EXT_API; static constexpr auto none = ROCPROFILER_HSA_FINALIZE_EXT_API_ID_NONE; static constexpr auto last = ROCPROFILER_HSA_FINALIZE_EXT_API_ID_LAST; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_HSA_FINALIZE_EXT_API; }; template <> @@ -101,6 +110,8 @@ struct hsa_domain_info static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_NONE; static constexpr auto none = 0; static constexpr auto last = 0; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_NONE; }; } // namespace hsa } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp index df241a2b90..5e001effab 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp @@ -26,6 +26,7 @@ #include #include +#include #include namespace rocprofiler @@ -125,6 +126,9 @@ struct hsa_api_impl static RetT functor(Args... args); }; +std::string_view +get_hsa_status_string(hsa_status_t _status); + template const char* name_by_id(uint32_t id); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp index 5dabf93ca7..c8fdc7f4a9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.cpp @@ -22,17 +22,24 @@ #include "lib/rocprofiler-sdk/hsa/queue.hpp" #include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/agent.hpp" #include "lib/rocprofiler-sdk/buffer.hpp" #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/hsa/code_object.hpp" #include "lib/rocprofiler-sdk/hsa/details/fmt.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" +#include "lib/rocprofiler-sdk/kernel_dispatch/tracing.hpp" +#include "lib/rocprofiler-sdk/registration.hpp" +#include "lib/rocprofiler-sdk/tracing/tracing.hpp" + +#include +#include +#include #include #include #include -#include #include #include @@ -69,93 +76,49 @@ namespace hsa { namespace { +template +inline bool +context_filter(const context::context* ctx, DomainT domain, Args... args) +{ + if constexpr(std::is_same::value) + { + return (ctx->buffered_tracer && ctx->buffered_tracer->domains(domain, args...)); + } + else if constexpr(std::is_same::value) + { + return (ctx->callback_tracer && ctx->callback_tracer->domains(domain, args...)); + } + else + { + static_assert(common::mpl::assert_false::value, "unsupported domain type"); + return false; + } +} + bool context_filter(const context::context* ctx) { - return (ctx->buffered_tracer && - (ctx->buffered_tracer->domains(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH))); + return (context_filter(ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH) || + context_filter(ctx, ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH)); } bool AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) { if(!data) return true; + + // if we have fully finalized, delete the data and return + if(registration::get_fini_status() > 0) + { + auto* _session = static_cast(data); + delete _session; + return false; + } + auto& queue_info_session = *static_cast(data); - // we need to decrement this reference count at the end of the functions - auto* _corr_id = queue_info_session.correlation_id; - // get the contexts that were active when the signal was created - const auto& ctxs = queue_info_session.contexts; - if(!ctxs.empty()) - { - // only do the following work if there are contexts that require this info - const auto* _rocp_agent = queue_info_session.rocp_agent; - auto _hsa_agent = queue_info_session.hsa_agent; - auto _queue_id = queue_info_session.queue_id; - auto _signal = queue_info_session.kernel_pkt.kernel_dispatch.completion_signal; - auto _kern_id = queue_info_session.kernel_id; - const auto& _extern_corr_ids = queue_info_session.extern_corr_ids; + kernel_dispatch::dispatch_complete(queue_info_session); - auto dispatch_time = hsa_amd_profiling_dispatch_time_t{}; - auto dispatch_time_status = - hsa::get_amd_ext_table()->hsa_amd_profiling_get_dispatch_time_fn( - _hsa_agent, _signal, &dispatch_time); - - // if we encounter this in CI, it will cause test to fail - ROCP_CI_LOG_IF( - ERROR, - dispatch_time_status == HSA_STATUS_SUCCESS && dispatch_time.end < dispatch_time.start) - << "hsa_amd_profiling_get_dispatch_time for kernel_id=" << _kern_id - << " on rocprofiler_agent=" << _rocp_agent->id.handle - << " returned dispatch times where the end time (" << dispatch_time.end - << ") was less than the start time (" << dispatch_time.start << ")"; - - for(const auto* itr : ctxs) - { - auto* _buffer = buffer::get_buffer( - itr->buffered_tracer->buffer_data.at(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)); - - // go ahead and create the correlation id value since we expect at least one of these - // domains will require it - auto _corr_id_v = - rocprofiler_correlation_id_t{.internal = 0, .external = context::null_user_data}; - if(_corr_id) - { - _corr_id_v.internal = _corr_id->internal; - _corr_id_v.external = _extern_corr_ids.at(itr); - } - - if(itr->buffered_tracer->domains(ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH)) - { - if(dispatch_time_status == HSA_STATUS_SUCCESS) - { - const auto& dispatch_packet = queue_info_session.kernel_pkt.kernel_dispatch; - - auto record = rocprofiler_buffer_tracing_kernel_dispatch_record_t{ - sizeof(rocprofiler_buffer_tracing_kernel_dispatch_record_t), - ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, - _corr_id_v, - dispatch_time.start, - dispatch_time.end, - _rocp_agent->id, - _queue_id, - _kern_id, - dispatch_packet.private_segment_size, - dispatch_packet.group_segment_size, - rocprofiler_dim3_t{dispatch_packet.workgroup_size_x, - dispatch_packet.workgroup_size_y, - dispatch_packet.workgroup_size_z}, - rocprofiler_dim3_t{dispatch_packet.grid_size_x, - dispatch_packet.grid_size_y, - dispatch_packet.grid_size_z}}; - - CHECK_NOTNULL(_buffer)->emplace(ROCPROFILER_BUFFER_CATEGORY_TRACING, - ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, - record); - } - } - } - } // Calls our internal callbacks to callers who need to be notified post // kernel execution. queue_info_session.queue.signal_callback([&](const auto& map) { @@ -186,6 +149,8 @@ AsyncSignalHandler(hsa_signal_value_t /*signal_v*/, void* data) queue_info_session.kernel_pkt.ext_amd_aql_pm4.completion_signal); } + // we need to decrement this reference count at the end of the functions + auto* _corr_id = queue_info_session.correlation_id; if(_corr_id) { LOG_IF(FATAL, _corr_id->get_ref_count() == 0) @@ -235,7 +200,11 @@ WriteInterceptor(const void* packets, void* data, hsa_amd_queue_intercept_packet_writer writer) { - using context_array_t = Queue::context_array_t; + using callback_record_t = Queue::queue_info_session_t::callback_record_t; + + // unique sequence id for the dispatch + static auto sequence_counter = std::atomic{0}; + auto&& CreateBarrierPacket = [](hsa_signal_t* dependency_signal, hsa_signal_t* completion_signal, std::vector& _packets) { @@ -249,33 +218,32 @@ WriteInterceptor(const void* packets, LOG_IF(FATAL, data == nullptr) << "WriteInterceptor was not passed a pointer to the queue"; - auto ctxs = context_array_t{}; - context::get_active_contexts(ctxs, context_filter); - auto& queue = *static_cast(data); // We have no packets or no one who needs to be notified, do nothing. - if(pkt_count == 0 || (queue.get_notifiers() == 0 && ctxs.empty())) + if(pkt_count == 0 || + (queue.get_notifiers() == 0 && context::get_active_contexts(context_filter).empty())) { writer(packets, pkt_count); return; } - auto thr_id = common::get_tid(); - auto* corr_id = context::get_latest_correlation_id(); - auto user_data = rocprofiler_user_data_t{.value = 0}; + auto tracing_data_v = tracing::tracing_data{}; + tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, + ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, + tracing_data_v); - // use thread-local value to reuse allocation - auto extern_corr_ids = Queue::queue_info_session_t::external_corr_id_map_t{}; + auto* corr_id = context::get_latest_correlation_id(); + auto thr_id = (corr_id) ? corr_id->thread_idx : common::get_tid(); + auto user_data = rocprofiler_user_data_t{.value = 0}; + auto internal_corr_id = (corr_id) ? corr_id->internal : 0; - // increase the reference count to denote that this correlation id is being used in a kernel - if(corr_id) - { - extern_corr_ids.clear(); // clear it so that it only contains the current contexts - extern_corr_ids.reserve(ctxs.size()); // reserve for performance - for(const auto* ctx : ctxs) - extern_corr_ids.emplace(ctx, ctx->correlation_tracer.external_correlator.get(thr_id)); - } + tracing::populate_external_correlation_ids( + tracing_data_v.external_correlation_ids, + thr_id, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH, + ROCPROFILER_KERNEL_DISPATCH_ENQUEUE, + internal_corr_id); const auto* packets_arr = static_cast(packets); auto transformed_packets = std::vector{}; @@ -300,18 +268,65 @@ WriteInterceptor(const void* packets, queue.create_signal(HSA_AMD_SIGNAL_AMD_GPU_ONLY, &kernel_pkt.ext_amd_aql_pm4.completion_signal); + // increase the reference count to denote that this correlation id is being used in a kernel + if(corr_id) + { + corr_id->add_ref_count(); + corr_id->add_kern_count(); + } + + auto dispatch_id = ++sequence_counter; + auto callback_record = + callback_record_t{sizeof(callback_record_t), + rocprofiler_timestamp_t{0}, + rocprofiler_timestamp_t{0}, + queue.get_agent().get_rocp_agent()->id, + queue.get_id(), + kernel_id, + dispatch_id, + kernel_pkt.kernel_dispatch.private_segment_size, + kernel_pkt.kernel_dispatch.group_segment_size, + rocprofiler_dim3_t{kernel_pkt.kernel_dispatch.workgroup_size_x, + kernel_pkt.kernel_dispatch.workgroup_size_y, + kernel_pkt.kernel_dispatch.workgroup_size_z}, + rocprofiler_dim3_t{kernel_pkt.kernel_dispatch.grid_size_x, + kernel_pkt.kernel_dispatch.grid_size_y, + kernel_pkt.kernel_dispatch.grid_size_z}}; + + { + auto tracer_data = callback_record; + tracing::execute_phase_enter_callbacks(tracing_data_v.callback_contexts, + thr_id, + internal_corr_id, + tracing_data_v.external_correlation_ids, + ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, + ROCPROFILER_KERNEL_DISPATCH_ENQUEUE, + tracer_data); + } + + // map all the external correlation ids (after enqueue enter phase) for all the contexts + // captured by the info session + tracing::update_external_correlation_ids( + tracing_data_v.external_correlation_ids, + thr_id, + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_KERNEL_DISPATCH); + // Stores the instrumentation pkt (i.e. AQL packets for counter collection) // along with an ID of the client we got the packet from (this will be returned via // completed_cb_t) - - inst_pkt_t inst_pkt; + auto inst_pkt = inst_pkt_t{}; // Signal callbacks that a kernel_pkt is being enqueued queue.signal_callback([&](const auto& map) { for(const auto& [client_id, cb_pair] : map) { - if(auto maybe_pkt = cb_pair.first( - queue, kernel_pkt, kernel_id, &user_data, extern_corr_ids, corr_id)) + if(auto maybe_pkt = cb_pair.first(queue, + kernel_pkt, + kernel_id, + dispatch_id, + &user_data, + tracing_data_v.external_correlation_ids, + corr_id)) { inst_pkt.push_back(std::make_pair(std::move(maybe_pkt), client_id)); } @@ -378,12 +393,6 @@ WriteInterceptor(const void* packets, LOG_IF(FATAL, packet_type != HSA_PACKET_TYPE_KERNEL_DISPATCH) << "get_kernel_id below might need to be updated"; - if(corr_id) - { - corr_id->add_ref_count(); - corr_id->add_kern_count(); - } - // Enqueue the signal into the handler. Will call completed_cb when // signal completes. queue.signal_async_handler( @@ -392,15 +401,20 @@ WriteInterceptor(const void* packets, .inst_pkt = std::move(inst_pkt), .interrupt_signal = interrupt_signal, .tid = thr_id, - .kernel_id = kernel_id, - .queue_id = queue.get_id(), .user_data = user_data, - .hsa_agent = queue.get_agent().get_hsa_agent(), - .rocp_agent = queue.get_agent().get_rocp_agent(), .correlation_id = corr_id, .kernel_pkt = kernel_pkt, - .contexts = ctxs, - .extern_corr_ids = extern_corr_ids}); + .callback_record = callback_record, + .tracing_data = tracing_data_v}); + + { + auto tracer_data = callback_record; + tracing::execute_phase_exit_callbacks(tracing_data_v.callback_contexts, + tracing_data_v.external_correlation_ids, + ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, + ROCPROFILER_KERNEL_DISPATCH_ENQUEUE, + tracer_data); + } } // Command is only executed if GLOG_v=2 or higher, otherwise it is a no-op diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp index 0144bd8bd9..ee2aa72c55 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp @@ -24,6 +24,7 @@ #include #include +#include #include #include "lib/common/container/small_vector.hpp" @@ -31,6 +32,8 @@ #include "lib/common/utility.hpp" #include "lib/rocprofiler-sdk/hsa/agent_cache.hpp" #include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" +#include "lib/rocprofiler-sdk/hsa/queue_info_session.hpp" +#include "lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp" #include #include @@ -40,6 +43,7 @@ #include #include +#include #include #include #include @@ -56,44 +60,6 @@ struct correlation_id; } // namespace context namespace hsa { -using ClientID = int64_t; -using inst_pkt_t = - common::container::small_vector, ClientID>, 4>; - -union rocprofiler_packet -{ - hsa_ext_amd_aql_pm4_packet_t ext_amd_aql_pm4; - hsa_kernel_dispatch_packet_t kernel_dispatch; - hsa_barrier_and_packet_t barrier_and; - hsa_barrier_or_packet_t barrier_or; - - rocprofiler_packet() - : ext_amd_aql_pm4{null_amd_aql_pm4_packet} - {} - - rocprofiler_packet(hsa_ext_amd_aql_pm4_packet_t val) - : ext_amd_aql_pm4{val} - {} - - rocprofiler_packet(hsa_kernel_dispatch_packet_t val) - : kernel_dispatch{val} - {} - - rocprofiler_packet(hsa_barrier_and_packet_t val) - : barrier_and{val} - {} - - rocprofiler_packet(hsa_barrier_or_packet_t val) - : barrier_or{val} - {} - - ~rocprofiler_packet() = default; - rocprofiler_packet(const rocprofiler_packet&) = default; - rocprofiler_packet(rocprofiler_packet&&) noexcept = default; - - rocprofiler_packet& operator=(const rocprofiler_packet&) = default; - rocprofiler_packet& operator=(rocprofiler_packet&&) noexcept = default; -}; enum class queue_state { normal = 0, @@ -105,31 +71,10 @@ enum class queue_state class Queue { public: - using context_t = context::context; - using context_array_t = common::container::small_vector; - using callback_t = void (*)(hsa_status_t status, hsa_queue_t* source, void* data); - - // Internal session information that is used by write interceptor - // to track state of the intercepted kernel. - struct queue_info_session_t - { - using external_corr_id_map_t = - std::unordered_map; - - Queue& queue; - inst_pkt_t inst_pkt = {}; - hsa_signal_t interrupt_signal = {}; - rocprofiler_thread_id_t tid = common::get_tid(); - rocprofiler_kernel_id_t kernel_id = 0; - rocprofiler_queue_id_t queue_id = {}; - rocprofiler_user_data_t user_data = {.value = 0}; - hsa_agent_t hsa_agent = {}; - const rocprofiler_agent_t* rocp_agent = nullptr; - context::correlation_id* correlation_id = nullptr; - rocprofiler_packet kernel_pkt = {}; - context_array_t contexts = {}; - external_corr_id_map_t extern_corr_ids = {}; - }; + using context_t = context::context; + using context_array_t = common::container::small_vector; + using callback_t = void (*)(hsa_status_t status, hsa_queue_t* source, void* data); + using queue_info_session_t = queue_info_session; // Function prototype used to notify consumers that a kernel has been // enqueued. An AQL packet can be returned that will be injected into @@ -137,7 +82,8 @@ public: using queue_cb_t = std::function( const Queue&, const rocprofiler_packet&, - uint64_t, + rocprofiler_kernel_id_t, + rocprofiler_dispatch_id_t, rocprofiler_user_data_t*, const queue_info_session_t::external_corr_id_map_t&, const context::correlation_id*)>; @@ -234,6 +180,5 @@ Queue::lock_queue(FuncT&& func) std::unique_lock lock(_lock_queue); func(); } - } // namespace hsa -} // namespace rocprofiler \ No newline at end of file +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp index a8a224a950..03c4ec0ecb 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_controller.cpp @@ -201,7 +201,7 @@ QueueController::init(CoreApiTable& core_table, AmdExtTable& ext_table) auto enable_intercepter = false; for(const auto& itr : context::get_registered_contexts()) { - constexpr auto expected_context_size = 160UL; + constexpr auto expected_context_size = 184UL; static_assert( sizeof(context::context) == expected_context_size + sizeof(std::shared_ptr), diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp new file mode 100644 index 0000000000..cb3ab287d3 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue_info_session.hpp @@ -0,0 +1,66 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include +#include + +#include "lib/common/container/small_vector.hpp" +#include "lib/common/utility.hpp" +#include "lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp" +#include "lib/rocprofiler-sdk/tracing/fwd.hpp" + +namespace rocprofiler +{ +namespace context +{ +struct context; +struct correlation_id; +} // namespace context + +namespace hsa +{ +class Queue; + +// Internal session information that is used by write interceptor +// to track state of the intercepted kernel. +struct queue_info_session +{ + using context_t = context::context; + using user_data_map_t = std::unordered_map; + using external_corr_id_map_t = user_data_map_t; + using callback_record_t = rocprofiler_callback_tracing_kernel_dispatch_data_t; + using context_array_t = common::container::small_vector; + + Queue& queue; + inst_pkt_t inst_pkt = {}; + hsa_signal_t interrupt_signal = {}; + rocprofiler_thread_id_t tid = common::get_tid(); + rocprofiler_user_data_t user_data = {.value = 0}; + context::correlation_id* correlation_id = nullptr; + rocprofiler_packet kernel_pkt = {}; + callback_record_t callback_record = {}; + tracing::tracing_data tracing_data = {}; +}; +} // namespace hsa +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp new file mode 100644 index 0000000000..ac0267a81a --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/rocprofiler_packet.hpp @@ -0,0 +1,77 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include +#include + +#include "lib/rocprofiler-sdk/hsa/aql_packet.hpp" + +#include +#include + +namespace rocprofiler +{ +namespace hsa +{ +using ClientID = int64_t; + +using inst_pkt_t = + common::container::small_vector, ClientID>, 4>; + +union rocprofiler_packet +{ + hsa_ext_amd_aql_pm4_packet_t ext_amd_aql_pm4; + hsa_kernel_dispatch_packet_t kernel_dispatch; + hsa_barrier_and_packet_t barrier_and; + hsa_barrier_or_packet_t barrier_or; + + rocprofiler_packet() + : ext_amd_aql_pm4{null_amd_aql_pm4_packet} + {} + + rocprofiler_packet(hsa_ext_amd_aql_pm4_packet_t val) + : ext_amd_aql_pm4{val} + {} + + rocprofiler_packet(hsa_kernel_dispatch_packet_t val) + : kernel_dispatch{val} + {} + + rocprofiler_packet(hsa_barrier_and_packet_t val) + : barrier_and{val} + {} + + rocprofiler_packet(hsa_barrier_or_packet_t val) + : barrier_or{val} + {} + + ~rocprofiler_packet() = default; + rocprofiler_packet(const rocprofiler_packet&) = default; + rocprofiler_packet(rocprofiler_packet&&) noexcept = default; + + rocprofiler_packet& operator=(const rocprofiler_packet&) = default; + rocprofiler_packet& operator=(rocprofiler_packet&&) noexcept = default; +}; +} // namespace hsa +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp index 867df07859..60bd00725f 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/scratch_memory.cpp @@ -28,8 +28,10 @@ #include "lib/rocprofiler-sdk/hsa/defines.hpp" #include "lib/rocprofiler-sdk/hsa/hsa.hpp" #include "lib/rocprofiler-sdk/hsa/queue_controller.hpp" +#include "lib/rocprofiler-sdk/tracing/tracing.hpp" #include +#include #include #include #include @@ -214,6 +216,8 @@ struct scratch_op_info using end_fn_t = scratch_op_info::function_t; \ static constexpr auto operation_idx = ROCPROFILER_##TOOL_OP; \ static constexpr auto name = #TOOL_OP; \ + static constexpr auto external_correlation_id_domain_idx = \ + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_SCRATCH_MEMORY; \ } SPECIALIZE_AMD_TOOL(SCRATCH_MEMORY_ALLOC, scratch_event_alloc_start, scratch_event_alloc_end); @@ -324,51 +328,6 @@ copy_table(hsa_amd_tool_table_t* _orig, uint64_t _tbl_instance) } } } - -struct buffered_context_data -{ - const context::context* ctx = nullptr; -}; - -struct callback_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_callback_tracing_record_t record = {}; - rocprofiler_user_data_t user_data = {.value = 0}; -}; - -void -populate_contexts(int operation_idx, - std::vector& callback_contexts, - std::vector& buffered_contexts) -{ - callback_contexts.clear(); - buffered_contexts.clear(); - - auto active_contexts = context::context_array_t{}; - for(const auto* itr : context::get_active_contexts(active_contexts)) - { - if(!itr) continue; - - if(itr->callback_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->callback_tracer->domains(ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, - operation_idx)) - callback_contexts.emplace_back( - callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); - } - - if(itr->buffered_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->buffered_tracer->domains(ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, - operation_idx)) - buffered_contexts.emplace_back(buffered_context_data{itr}); - } - } -} - } // namespace static_assert(ROCPROFILER_SCRATCH_MEMORY_ALLOC == @@ -442,8 +401,9 @@ get_tls_pair(rocprofiler_callback_phase_t phase) { callback_data_t callback_data = common::init_public_api_struct(callback_data_t{}); buffered_data_t buffered_data = common::init_public_api_struct(buffered_data_t{}); - std::vector callback_ctxs = {}; - std::vector buffered_ctxs = {}; + tracing::callback_context_data_vec_t callback_contexts = {}; + tracing::buffered_context_data_vec_t buffered_contexts = {}; + tracing::external_correlation_id_map_t external_corr_ids = {}; }; static thread_local auto tls = tls_data{}; @@ -451,9 +411,19 @@ get_tls_pair(rocprofiler_callback_phase_t phase) if(phase == ROCPROFILER_CALLBACK_PHASE_ENTER) { + // since the context data structures are TLS, we need to clear the contexts first to prevent + // duplicate entries + using clear_containers = std::true_type; + LOG_IF(FATAL, held) << "Overwriting scratch memory TLS data"; held = true; - populate_contexts(OpIdx, tls.callback_ctxs, tls.buffered_ctxs); + tracing::populate_contexts(ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, + ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, + OpIdx, + tls.callback_contexts, + tls.buffered_contexts, + tls.external_corr_ids, + clear_containers{}); } else { @@ -473,6 +443,8 @@ impl(Args... args) constexpr auto OpIdx = scratch_op_info::operation; constexpr auto OpPhase = scratch_op_info::phase; + constexpr auto external_corr_id_domain_idx = + amd_tool_api_info::external_correlation_id_domain_idx; auto&& _tied_args = std::tie(args...); auto& event_data = std::get<0>(_tied_args); @@ -480,9 +452,11 @@ impl(Args... args) // this lets start and end of the same type have the same thread local storage auto& tls = get_tls_pair(OpPhase); - if(tls.callback_ctxs.empty() && tls.buffered_ctxs.empty()) return HSA_STATUS_SUCCESS; + if(tls.callback_contexts.empty() && tls.buffered_contexts.empty()) return HSA_STATUS_SUCCESS; - const auto tid = common::get_tid(); + const auto thr_id = common::get_tid(); + auto* corr_id = context::get_latest_correlation_id(); + auto internal_corr_id = (corr_id) ? corr_id->internal : 0; [[maybe_unused]] const auto get_agent_id = [](const hsa_queue_t* hsa_queue) -> rocprofiler_agent_id_t { @@ -503,37 +477,9 @@ impl(Args... args) return _agent_id; }; - auto* corr_id = context::get_latest_correlation_id(); - const auto invoke_callbacks = [&](const rocprofiler_callback_phase_t _phase) { - for(auto& itr : tls.callback_ctxs) - { - auto corr_id_v = rocprofiler_correlation_id_t{}; - if(corr_id) - { - corr_id_v.internal = corr_id->internal; - corr_id_v.external = itr.ctx->correlation_tracer.external_correlator.get(tid); - } - - auto record = rocprofiler_callback_tracing_record_t{ - .context_id = rocprofiler_context_id_t{itr.ctx->context_idx}, - .thread_id = tid, - .correlation_id = corr_id_v, - .kind = ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, - .operation = OpIdx, - .phase = _phase, - .payload = static_cast(&tls.callback_data), - }; - - auto& cb_info = itr.ctx->callback_tracer->callback_data.at( - ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY); - - cb_info.callback(record, &itr.user_data, &cb_info.data); - } - }; - if constexpr(OpPhase == ROCPROFILER_CALLBACK_PHASE_ENTER) { - if(!tls.callback_ctxs.empty()) + if(!tls.callback_contexts.empty()) { tls.callback_data.agent_id = get_agent_id(event_data.scratch_alloc_start->queue); tls.callback_data.queue_id = {event_data.scratch_alloc_start->queue->id}; @@ -546,28 +492,44 @@ impl(Args... args) event_data.scratch_alloc_start->dispatch_id; } - invoke_callbacks(OpPhase); // NOLINT readability-misleading-indentation + tracing::populate_external_correlation_ids(tls.external_corr_ids, + thr_id, + external_corr_id_domain_idx, + OpIdx, + internal_corr_id); + + tracing::execute_phase_enter_callbacks(tls.callback_contexts, + thr_id, + internal_corr_id, + tls.external_corr_ids, + ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, + OpIdx, + tls.callback_data); + + // enter callback may update the external correlation id field + tracing::update_external_correlation_ids( + tls.external_corr_ids, thr_id, external_corr_id_domain_idx); } - if(!tls.buffered_ctxs.empty()) + if(!tls.buffered_contexts.empty()) { tls.buffered_data.kind = ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY; tls.buffered_data.operation = OpIdx; tls.buffered_data.agent_id = get_agent_id(event_data.scratch_alloc_start->queue); tls.buffered_data.queue_id = {event_data.scratch_alloc_start->queue->id}; - tls.buffered_data.thread_id = tid; + tls.buffered_data.thread_id = thr_id; tls.buffered_data.start_timestamp = common::timestamp_ns(); } } else if constexpr(OpPhase == ROCPROFILER_CALLBACK_PHASE_EXIT) { - if(!tls.buffered_ctxs.empty()) + if(!tls.buffered_contexts.empty()) { tls.buffered_data.flags = get_flags(event_data); tls.buffered_data.end_timestamp = common::timestamp_ns(); } - if(!tls.callback_ctxs.empty()) + if(!tls.callback_contexts.empty()) { tls.callback_data.flags = get_flags(event_data); tls.callback_data.args_kind = event_data.none->kind; @@ -579,31 +541,23 @@ impl(Args... args) data_args.num_slots = event_data.scratch_alloc_end->num_slots; } - invoke_callbacks(OpPhase); // NOLINT readability-misleading-indentation + tracing::execute_phase_exit_callbacks(tls.callback_contexts, + tls.external_corr_ids, + ROCPROFILER_CALLBACK_TRACING_SCRATCH_MEMORY, + OpIdx, + tls.callback_data); } - if(!tls.buffered_ctxs.empty()) + if(!tls.buffered_contexts.empty()) { - for(const auto& itr : tls.buffered_ctxs) - { - auto* _buffer = buffer::get_buffer(itr.ctx->buffered_tracer->buffer_data.at( - ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY)); - - auto corr_id_v = rocprofiler_correlation_id_t{}; - if(corr_id) - { - // TODO(mkuriche) should the id be generated at entry? - corr_id_v.internal = corr_id->internal; - corr_id_v.external = itr.ctx->correlation_tracer.external_correlator.get(tid); - } - - auto _record = tls.buffered_data; - _record.correlation_id = corr_id_v; - - CHECK_NOTNULL(_buffer)->emplace(ROCPROFILER_BUFFER_CATEGORY_TRACING, - ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, - _record); - } + auto _buffered_data = tls.buffered_data; + tracing::execute_buffer_record_emplace(tls.buffered_contexts, + thr_id, + internal_corr_id, + tls.external_corr_ids, + ROCPROFILER_BUFFER_TRACING_SCRATCH_MEMORY, + OpIdx, + std::move(_buffered_data)); } } diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/CMakeLists.txt new file mode 100644 index 0000000000..a322b5ab17 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/CMakeLists.txt @@ -0,0 +1,7 @@ +# +set(ROCPROFILER_LIB_KERNEL_DISPATCH_SOURCES kernel_dispatch.cpp tracing.cpp) +set(ROCPROFILER_LIB_KERNEL_DISPATCH_HEADERS kernel_dispatch.hpp tracing.hpp) + +target_sources( + rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_KERNEL_DISPATCH_SOURCES} + ${ROCPROFILER_LIB_KERNEL_DISPATCH_HEADERS}) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.cpp new file mode 100644 index 0000000000..c9bbedf01e --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.cpp @@ -0,0 +1,132 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include "lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp" + +#include + +#if defined(ROCPROFILER_CI) +# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__) +# define ROCP_CI_LOG(NON_CI_LEVEL, ...) ROCP_FATAL +#else +# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__) +# define ROCP_CI_LOG(NON_CI_LEVEL, ...) LOG(NON_CI_LEVEL) +#endif + +namespace rocprofiler +{ +namespace kernel_dispatch +{ +namespace +{ +#define ROCPROFILER_KERNEL_DISPATCH_INFO(CODE) \ + template <> \ + struct kernel_dispatch_info \ + { \ + static constexpr auto operation_idx = ROCPROFILER_KERNEL_DISPATCH_##CODE; \ + static constexpr auto name = #CODE; \ + }; + +template +struct kernel_dispatch_info; + +ROCPROFILER_KERNEL_DISPATCH_INFO(NONE) +ROCPROFILER_KERNEL_DISPATCH_INFO(ENQUEUE) +ROCPROFILER_KERNEL_DISPATCH_INFO(COMPLETE) + +template +const char* +name_by_id(const uint32_t id, std::index_sequence) +{ + if(Idx == id) return kernel_dispatch_info::name; + if constexpr(sizeof...(IdxTail) > 0) + return name_by_id(id, std::index_sequence{}); + else + return nullptr; +} + +template +uint32_t +id_by_name(const char* name, std::index_sequence) +{ + if(std::string_view{kernel_dispatch_info::name} == std::string_view{name}) + return kernel_dispatch_info::operation_idx; + if constexpr(sizeof...(IdxTail) > 0) + return id_by_name(name, std::index_sequence{}); + else + return ROCPROFILER_HSA_AMD_EXT_API_ID_NONE; +} + +template +void +get_ids(std::vector& _id_list, std::index_sequence) +{ + auto _emplace = [](auto& _vec, uint32_t _v) { + if(_v < static_cast(ROCPROFILER_HSA_AMD_EXT_API_ID_LAST)) _vec.emplace_back(_v); + }; + + (_emplace(_id_list, kernel_dispatch_info::operation_idx), ...); +} + +template +void +get_names(std::vector& _name_list, std::index_sequence) +{ + auto _emplace = [](auto& _vec, const char* _v) { + if(_v != nullptr && strnlen(_v, 1) > 0) _vec.emplace_back(_v); + }; + + (_emplace(_name_list, kernel_dispatch_info::name), ...); +} +} // namespace + +const char* +name_by_id(uint32_t id) +{ + return name_by_id(id, std::make_index_sequence{}); +} + +uint32_t +id_by_name(const char* name) +{ + return id_by_name(name, std::make_index_sequence{}); +} + +std::vector +get_ids() +{ + auto _data = std::vector{}; + _data.reserve(ROCPROFILER_KERNEL_DISPATCH_LAST); + get_ids(_data, std::make_index_sequence{}); + return _data; +} + +std::vector +get_names() +{ + auto _data = std::vector{}; + _data.reserve(ROCPROFILER_KERNEL_DISPATCH_LAST); + get_names(_data, std::make_index_sequence{}); + return _data; +} +} // namespace kernel_dispatch +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp new file mode 100644 index 0000000000..887cf5ba0f --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/kernel_dispatch.hpp @@ -0,0 +1,48 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include "lib/rocprofiler-sdk/hsa/queue_info_session.hpp" + +#include + +#include +#include + +namespace rocprofiler +{ +namespace kernel_dispatch +{ +const char* +name_by_id(uint32_t id); + +uint32_t +id_by_name(const char* name); + +std::vector +get_names(); + +std::vector +get_ids(); +} // namespace kernel_dispatch +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp new file mode 100644 index 0000000000..e443177210 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.cpp @@ -0,0 +1,145 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#include "lib/rocprofiler-sdk/kernel_dispatch/tracing.hpp" +#include "lib/rocprofiler-sdk/agent.hpp" +#include "lib/rocprofiler-sdk/buffer.hpp" +#include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/hsa/hsa.hpp" +#include "lib/rocprofiler-sdk/hsa/queue.hpp" +#include "lib/rocprofiler-sdk/tracing/tracing.hpp" + +#include +#include + +#include + +#include + +#if defined(ROCPROFILER_CI) +# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(FATAL, __VA_ARGS__) +# define ROCP_CI_LOG(NON_CI_LEVEL, ...) ROCP_FATAL +#else +# define ROCP_CI_LOG_IF(NON_CI_LEVEL, ...) LOG_IF(NON_CI_LEVEL, __VA_ARGS__) +# define ROCP_CI_LOG(NON_CI_LEVEL, ...) LOG(NON_CI_LEVEL) +#endif + +namespace rocprofiler +{ +namespace kernel_dispatch +{ +namespace +{ +using queue_info_session_t = hsa::queue_info_session; +using kernel_dispatch_record_t = rocprofiler_buffer_tracing_kernel_dispatch_record_t; +} // namespace + +void +dispatch_complete(queue_info_session_t& session) +{ + // get the contexts that were active when the signal was created + auto& tracing_data_v = session.tracing_data; + if(tracing_data_v.callback_contexts.empty() && tracing_data_v.buffered_contexts.empty()) return; + + // we need to decrement this reference count at the end of the functions + auto* _corr_id = session.correlation_id; + + // only do the following work if there are contexts that require this info + auto& callback_record = session.callback_record; + const auto& _extern_corr_ids = session.tracing_data.external_correlation_ids; + const auto* _rocp_agent = agent::get_agent(callback_record.agent_id); + auto _hsa_agent = agent::get_hsa_agent(_rocp_agent); + auto _kern_id = callback_record.kernel_id; + auto _signal = session.kernel_pkt.kernel_dispatch.completion_signal; + auto _tid = session.tid; + + auto dispatch_time = hsa_amd_profiling_dispatch_time_t{}; + auto dispatch_time_status = + (_hsa_agent) ? hsa::get_amd_ext_table()->hsa_amd_profiling_get_dispatch_time_fn( + *_hsa_agent, _signal, &dispatch_time) + : HSA_STATUS_ERROR; + + if(dispatch_time_status == HSA_STATUS_SUCCESS) + { + callback_record.start_timestamp = dispatch_time.start; + callback_record.end_timestamp = dispatch_time.end; + } + + // if we encounter this in CI, it will cause test to fail + ROCP_CI_LOG_IF( + ERROR, + dispatch_time_status == HSA_STATUS_SUCCESS && dispatch_time.end < dispatch_time.start) + << "hsa_amd_profiling_get_dispatch_time for kernel_id=" << _kern_id + << " on rocprofiler_agent=" << _rocp_agent->id.handle + << " returned dispatch times where the end time (" << dispatch_time.end + << ") was less than the start time (" << dispatch_time.start << ")"; + + ROCP_CI_LOG_IF(ERROR, dispatch_time_status != HSA_STATUS_SUCCESS) + << "hsa_amd_profiling_get_dispatch_time for kernel id=" << _kern_id << " returned " + << dispatch_time_status << " :: " << hsa::get_hsa_status_string(dispatch_time_status); + + auto _internal_corr_id = (_corr_id) ? _corr_id->internal : 0; + + if(dispatch_time_status == HSA_STATUS_SUCCESS) + { + if(!tracing_data_v.callback_contexts.empty()) + { + auto tracer_data = callback_record; + tracing::execute_phase_none_callbacks(tracing_data_v.callback_contexts, + _tid, + _internal_corr_id, + _extern_corr_ids, + ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, + ROCPROFILER_KERNEL_DISPATCH_COMPLETE, + tracer_data); + } + + if(!tracing_data_v.buffered_contexts.empty()) + { + auto record = kernel_dispatch_record_t{sizeof(kernel_dispatch_record_t), + ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, + ROCPROFILER_KERNEL_DISPATCH_COMPLETE, + rocprofiler_correlation_id_t{}, + _tid, + callback_record.start_timestamp, + callback_record.end_timestamp, + callback_record.agent_id, + callback_record.queue_id, + callback_record.kernel_id, + callback_record.dispatch_id, + callback_record.private_segment_size, + callback_record.group_segment_size, + callback_record.workgroup_size, + callback_record.grid_size}; + + tracing::execute_buffer_record_emplace(tracing_data_v.buffered_contexts, + _tid, + _internal_corr_id, + _extern_corr_ids, + ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, + ROCPROFILER_KERNEL_DISPATCH_COMPLETE, + record); + } + } +} +} // namespace kernel_dispatch +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.hpp new file mode 100644 index 0000000000..0a15672577 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/kernel_dispatch/tracing.hpp @@ -0,0 +1,39 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/hsa/queue_info_session.hpp" + +namespace rocprofiler +{ +namespace kernel_dispatch +{ +using context_t = context::context; +using user_data_map_t = std::unordered_map; +using external_corr_id_map_t = user_data_map_t; + +void +dispatch_complete(hsa::queue_info_session&); +} // namespace kernel_dispatch +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp index 63e36e44a2..ff4710003c 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.cpp @@ -28,6 +28,7 @@ #include "lib/rocprofiler-sdk/context/context.hpp" #include "lib/rocprofiler-sdk/marker/utils.hpp" #include "lib/rocprofiler-sdk/registration.hpp" +#include "lib/rocprofiler-sdk/tracing/tracing.hpp" #include #include @@ -130,57 +131,6 @@ roctx_api_impl::exec(FuncT&& _func, Args&&... args) return get_default_retval(); } -namespace -{ -using correlation_service = context::correlation_tracing_service; - -struct callback_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_callback_tracing_record_t record = {}; - rocprofiler_user_data_t user_data = {.value = 0}; -}; - -struct buffered_context_data -{ - const context::context* ctx = nullptr; - rocprofiler_user_data_t external_correlation = {}; -}; - -constexpr auto empty_user_data = rocprofiler_user_data_t{.value = 0}; - -void -populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, - rocprofiler_buffer_tracing_kind_t buffered_domain_idx, - int operation_idx, - std::vector& callback_contexts, - std::vector& buffered_contexts) -{ - auto active_contexts = context::context_array_t{}; - auto thr_id = common::get_tid(); - for(const auto* itr : context::get_active_contexts(active_contexts)) - { - if(!itr) continue; - - if(itr->callback_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->callback_tracer->domains(callback_domain_idx, operation_idx)) - callback_contexts.emplace_back( - callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); - } - - if(itr->buffered_tracer) - { - // if the given domain + op is not enabled, skip this context - if(itr->buffered_tracer->domains(buffered_domain_idx, operation_idx)) - buffered_contexts.emplace_back(buffered_context_data{ - itr, itr->correlation_tracer.external_correlator.get(thr_id)}); - } - } -} -} // namespace - template template RetT @@ -190,17 +140,22 @@ roctx_api_impl::functor(Args... args) using callback_api_data_t = typename roctx_domain_info::callback_data_type; using buffered_api_data_t = typename roctx_domain_info::buffer_data_type; + constexpr auto external_corr_id_domain_idx = + roctx_domain_info::external_correlation_id_domain_idx; + LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name; auto thr_id = common::get_tid(); - auto callback_contexts = std::vector{}; - auto buffered_contexts = std::vector{}; + auto callback_contexts = tracing::callback_context_data_vec_t{}; + auto buffered_contexts = tracing::buffered_context_data_vec_t{}; + auto external_corr_ids = tracing::external_correlation_id_map_t{}; - populate_contexts(info_type::callback_domain_idx, - info_type::buffered_domain_idx, - info_type::operation_idx, - callback_contexts, - buffered_contexts); + tracing::populate_contexts(info_type::callback_domain_idx, + info_type::buffered_domain_idx, + info_type::operation_idx, + callback_contexts, + buffered_contexts, + external_corr_ids); if(callback_contexts.empty() && buffered_contexts.empty()) { @@ -211,68 +166,42 @@ roctx_api_impl::functor(Args... args) return; } + LOG_IF(FATAL, external_corr_ids.size() < (callback_contexts.size() + buffered_contexts.size())) + << "missing external correlation ids"; + auto ref_count = 2; auto buffer_record = common::init_public_api_struct(buffered_api_data_t{}); auto tracer_data = common::init_public_api_struct(callback_api_data_t{}); - auto* corr_id = correlation_service::construct(ref_count); + auto* corr_id = tracing::correlation_service::construct(ref_count); auto internal_corr_id = corr_id->internal; - // construct the buffered info before the callback so the callbacks are as closely wrapped - // around the function call as possible - if(!buffered_contexts.empty()) - { - buffer_record.kind = info_type::buffered_domain_idx; - // external correlation will be updated right before record is placed in buffer - buffer_record.correlation_id = - rocprofiler_correlation_id_t{internal_corr_id, empty_user_data}; - buffer_record.operation = info_type::operation_idx; - buffer_record.thread_id = thr_id; - } - - tracer_data.size = sizeof(callback_api_data_t); - set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward(args)...); + tracing::populate_external_correlation_ids(external_corr_ids, + thr_id, + external_corr_id_domain_idx, + info_type::operation_idx, + internal_corr_id); // invoke the callbacks if(!callback_contexts.empty()) { set_data_args(info_type::get_api_data_args(tracer_data.args), std::forward(args)...); - for(auto& itr : callback_contexts) - { - auto& ctx = itr.ctx; - auto& record = itr.record; - auto& user_data = itr.user_data; - - auto extern_corr_id_v = ctx->correlation_tracer.external_correlator.get(thr_id); - - auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; - record = - rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, - thr_id, - corr_id_v, - info_type::callback_domain_idx, - info_type::operation_idx, - ROCPROFILER_CALLBACK_PHASE_ENTER, - static_cast(&tracer_data)}; - - auto& callback_info = - ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); - callback_info.callback(record, &user_data, callback_info.data); - - // enter callback may update the external correlation id field - record.correlation_id.external = - ctx->correlation_tracer.external_correlator.get(thr_id); - } + tracing::execute_phase_enter_callbacks(callback_contexts, + thr_id, + internal_corr_id, + external_corr_ids, + info_type::callback_domain_idx, + info_type::operation_idx, + tracer_data); } + // enter callback may update the external correlation id field + tracing::update_external_correlation_ids( + external_corr_ids, thr_id, external_corr_id_domain_idx); + // record the start timestamp as close to the function call as possible if(!buffered_contexts.empty()) { - for(auto& itr : buffered_contexts) - { - itr.external_correlation = itr.ctx->correlation_tracer.external_correlator.get(thr_id); - } - buffer_record.start_timestamp = common::timestamp_ns(); } @@ -291,41 +220,22 @@ roctx_api_impl::functor(Args... args) { set_data_retval(tracer_data.retval, _ret); - for(auto& itr : callback_contexts) - { - auto& ctx = itr.ctx; - auto& record = itr.record; - auto& user_data = itr.user_data; - - record.phase = ROCPROFILER_CALLBACK_PHASE_EXIT; - record.payload = static_cast(&tracer_data); - - auto& callback_info = - ctx->callback_tracer->callback_data.at(info_type::callback_domain_idx); - callback_info.callback(record, &user_data, callback_info.data); - } + tracing::execute_phase_exit_callbacks(callback_contexts, + external_corr_ids, + info_type::callback_domain_idx, + info_type::operation_idx, + tracer_data); } if(!buffered_contexts.empty()) { - for(auto& itr : buffered_contexts) - { - assert(itr.ctx->buffered_tracer); - auto buffer_id = - itr.ctx->buffered_tracer->buffer_data.at(info_type::buffered_domain_idx); - auto buffer_v = buffer::get_buffer(buffer_id); - if(buffer_v && buffer_v->context_id == itr.ctx->context_idx && - buffer_v->buffer_id == buffer_id.handle) - { - // make copy of record - auto record_v = buffer_record; - // update the record with the correlation - record_v.correlation_id.external = itr.external_correlation; - - buffer_v->emplace( - ROCPROFILER_BUFFER_CATEGORY_TRACING, info_type::buffered_domain_idx, record_v); - } - } + tracing::execute_buffer_record_emplace(buffered_contexts, + thr_id, + internal_corr_id, + external_corr_ids, + info_type::buffered_domain_idx, + info_type::operation_idx, + buffer_record); } // decrement the reference count after usage in the callback/buffers diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.def.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.def.cpp index 50e9e9003b..32003db8d8 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.def.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/marker/marker.def.cpp @@ -23,6 +23,7 @@ #include "lib/rocprofiler-sdk/marker/defines.hpp" #include "lib/rocprofiler-sdk/marker/marker.hpp" +#include #include namespace rocprofiler @@ -42,33 +43,39 @@ template <> struct roctx_domain_info : roctx_domain_info { + using enum_type = rocprofiler_marker_core_api_id_t; static constexpr auto callback_domain_idx = ROCPROFILER_CALLBACK_TRACING_MARKER_CORE_API; static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API; static constexpr auto none = ROCPROFILER_MARKER_CORE_API_ID_NONE; static constexpr auto last = ROCPROFILER_MARKER_CORE_API_ID_LAST; - using enum_type = rocprofiler_marker_core_api_id_t; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MARKER_CORE_API; }; template <> struct roctx_domain_info : roctx_domain_info { + using enum_type = rocprofiler_marker_control_api_id_t; static constexpr auto callback_domain_idx = ROCPROFILER_CALLBACK_TRACING_MARKER_CONTROL_API; static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_MARKER_CONTROL_API; static constexpr auto none = ROCPROFILER_MARKER_CONTROL_API_ID_NONE; static constexpr auto last = ROCPROFILER_MARKER_CONTROL_API_ID_LAST; - using enum_type = rocprofiler_marker_control_api_id_t; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MARKER_CONTROL_API; }; template <> struct roctx_domain_info : roctx_domain_info { + using enum_type = rocprofiler_marker_name_api_id_t; static constexpr auto callback_domain_idx = ROCPROFILER_CALLBACK_TRACING_MARKER_NAME_API; static constexpr auto buffered_domain_idx = ROCPROFILER_BUFFER_TRACING_MARKER_NAME_API; static constexpr auto none = ROCPROFILER_MARKER_NAME_API_ID_NONE; static constexpr auto last = ROCPROFILER_MARKER_NAME_API_ID_LAST; - using enum_type = rocprofiler_marker_name_api_id_t; + static constexpr auto external_correlation_id_domain_idx = + ROCPROFILER_EXTERNAL_CORRELATION_REQUEST_MARKER_NAME_API; }; } // namespace marker } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/CMakeLists.txt index 5828bfcc0f..73d44e08e2 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tests/CMakeLists.txt @@ -28,7 +28,7 @@ gtest_add_tests( TEST_LIST lib_TESTS WORKING_DIRECTORY ${CMAKE_CURRENT_BINARY_DIR}) -set_tests_properties(${lib_TESTS} PROPERTIES TIMEOUT 45 LABELS "unittests") +set_tests_properties(${lib_TESTS} PROPERTIES TIMEOUT 30 LABELS "unittests") # -------------------------------------------------------------------------------------- # # @@ -62,5 +62,5 @@ set(rocprofiler-lib-tests-env "LD_LIBRARY_PATH=${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}:$ENV{LD_LIBRARY_PATH}") set_tests_properties( - ${shared_lib_TESTS} PROPERTIES TIMEOUT 360 LABELS "unittests" ENVIRONMENT + ${shared_lib_TESTS} PROPERTIES TIMEOUT 120 LABELS "unittests" ENVIRONMENT "${rocprofiler-lib-tests-env}") diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp index 3ba762425a..0a645ce081 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/thread_trace/att_core.cpp @@ -169,12 +169,14 @@ ThreadTracer::start_context() std::nullopt, [=](const hsa::Queue& q, const hsa::rocprofiler_packet& kern_pkt, - uint64_t kernel_id, + rocprofiler_kernel_id_t kernel_id, + rocprofiler_dispatch_id_t dispatch_id, rocprofiler_user_data_t* user_data, const hsa::Queue::queue_info_session_t::external_corr_id_map_t& extern_corr_ids, const context::correlation_id* corr_id) { - (void) user_data; return pre_kernel_call(*this, q, kern_pkt, kernel_id, extern_corr_ids, corr_id); + (void) user_data; + (void) dispatch_id; }, [=](const hsa::Queue& q, hsa::rocprofiler_packet kern_pkt, diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/CMakeLists.txt new file mode 100644 index 0000000000..16a9dd487a --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/CMakeLists.txt @@ -0,0 +1,6 @@ +# +set(ROCPROFILER_LIB_TRACING_SOURCES) +set(ROCPROFILER_LIB_TRACING_HEADERS fwd.hpp tracing.hpp) + +target_sources(rocprofiler-object-library PRIVATE ${ROCPROFILER_LIB_TRACING_SOURCES} + ${ROCPROFILER_LIB_TRACING_HEADERS}) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/fwd.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/fwd.hpp new file mode 100644 index 0000000000..eb7f40a4e9 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/fwd.hpp @@ -0,0 +1,74 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include "lib/common/container/small_vector.hpp" + +#include + +#include +#include +#include + +namespace rocprofiler +{ +namespace context +{ +struct context; +struct correlation_tracing_service; +} // namespace context +namespace tracing +{ +template +using small_vector_t = common::container::small_vector; +using correlation_service = context::correlation_tracing_service; +using context_t = context::context; +using context_array_t = common::container::small_vector; +using external_correlation_id_map_t = std::unordered_map; + +constexpr auto context_data_vec_size = 2; +constexpr auto empty_user_data = rocprofiler_user_data_t{.value = 0}; + +struct callback_context_data +{ + const context_t* ctx = nullptr; + rocprofiler_callback_tracing_record_t record = {}; + rocprofiler_user_data_t user_data = {.value = 0}; +}; + +struct buffered_context_data +{ + const context_t* ctx = nullptr; +}; + +using callback_context_data_vec_t = small_vector_t; +using buffered_context_data_vec_t = small_vector_t; + +struct tracing_data +{ + callback_context_data_vec_t callback_contexts = {}; + buffered_context_data_vec_t buffered_contexts = {}; + external_correlation_id_map_t external_correlation_ids = {}; +}; +} // namespace tracing +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp new file mode 100644 index 0000000000..49f8863ab7 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/tracing/tracing.hpp @@ -0,0 +1,420 @@ +// MIT License +// +// Copyright (c) 2023 Advanced Micro Devices, Inc. All rights reserved. +// +// Permission is hereby granted, free of charge, to any person obtaining a copy +// of this software and associated documentation files (the "Software"), to deal +// in the Software without restriction, including without limitation the rights +// to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +// copies of the Software, and to permit persons to whom the Software is +// furnished to do so, subject to the following conditions: +// +// The above copyright notice and this permission notice shall be included in +// all copies or substantial portions of the Software. +// +// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +// THE SOFTWARE. + +#pragma once + +#include "lib/common/mpl.hpp" +#include "lib/rocprofiler-sdk/buffer.hpp" +#include "lib/rocprofiler-sdk/context/context.hpp" +#include "lib/rocprofiler-sdk/tracing/fwd.hpp" + +#include + +#include +#include +#include + +namespace rocprofiler +{ +namespace tracing +{ +// template +// bool +// context_filter(const context::context* ctx, DomainT domain, Args... args); + +// template +// void +// populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, +// rocprofiler_buffer_tracing_kind_t buffered_domain_idx, +// rocprofiler_tracing_operation_t operation_idx, +// callback_context_data_vec_t& callback_contexts, +// buffered_context_data_vec_t& buffered_contexts, +// external_correlation_id_map_t& extern_corr_ids, +// ClearContainersT = ClearContainersT{}); + +// template +// void +// populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, +// rocprofiler_buffer_tracing_kind_t buffered_domain_idx, +// callback_context_data_vec_t& callback_contexts, +// buffered_context_data_vec_t& buffered_contexts, +// external_correlation_id_map_t& extern_corr_ids, +// ClearContainersT = ClearContainersT{}); + +// template +// void +// populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, +// rocprofiler_buffer_tracing_kind_t buffered_domain_idx, +// rocprofiler_tracing_operation_t operation_idx, +// tracing_data& data, +// ClearContainersT = ClearContainersT{}); + +// template +// void +// populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, +// rocprofiler_buffer_tracing_kind_t buffered_domain_idx, +// tracing_data& data, +// ClearContainersT = ClearContainersT{}); + +// void +// populate_external_correlation_ids(external_correlation_id_map_t& external_corr_ids, +// rocprofiler_thread_id_t thr_id, +// rocprofiler_external_correlation_id_request_kind_t kind, +// rocprofiler_tracing_operation_t operation, +// uint64_t internal_corr_id); + +// void +// update_external_correlation_ids(external_correlation_id_map_t& external_corr_ids, +// rocprofiler_thread_id_t thr_id, +// rocprofiler_external_correlation_id_request_kind_t kind); + +// template +// void +// execute_phase_none_callbacks(callback_context_data_vec_t& callback_contexts, +// rocprofiler_thread_id_t thr_id, +// uint64_t internal_corr_id, +// external_correlation_id_map_t& external_corr_ids, +// rocprofiler_callback_tracing_kind_t domain, +// rocprofiler_tracing_operation_t operation, +// TracerDataT& tracer_data); + +// template +// void +// execute_phase_enter_callbacks(callback_context_data_vec_t& callback_contexts, +// rocprofiler_thread_id_t thr_id, +// uint64_t internal_corr_id, +// external_correlation_id_map_t& external_corr_ids, +// rocprofiler_callback_tracing_kind_t domain, +// rocprofiler_tracing_operation_t operation, +// TracerDataT& tracer_data); + +// template +// void +// execute_phase_exit_callbacks(callback_context_data_vec_t& callback_contexts, +// external_correlation_id_map_t& external_corr_ids, +// rocprofiler_callback_tracing_kind_t domain, +// rocprofiler_tracing_operation_t operation, +// TracerDataT& tracer_data); + +// template +// void +// execute_buffer_record_emplace(buffered_context_data_vec_t& buffered_contexts, +// rocprofiler_thread_id_t thr_id, +// uint64_t internal_corr_id, +// external_correlation_id_map_t& external_corr_ids, +// rocprofiler_buffer_tracing_kind_t domain, +// OperationT operation, +// BufferRecordT&& base_record); + +template +inline bool +context_filter(const context::context* ctx, DomainT domain, Args... args) +{ + if constexpr(std::is_same::value) + { + return (ctx->buffered_tracer && ctx->buffered_tracer->domains(domain, args...)); + } + else if constexpr(std::is_same::value) + { + return (ctx->callback_tracer && ctx->callback_tracer->domains(domain, args...)); + } + else + { + static_assert(common::mpl::assert_false::value, "unsupported domain type"); + return false; + } +} + +template +inline void +populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, + rocprofiler_buffer_tracing_kind_t buffered_domain_idx, + rocprofiler_tracing_operation_t operation_idx, + callback_context_data_vec_t& callback_contexts, + buffered_context_data_vec_t& buffered_contexts, + external_correlation_id_map_t& extern_corr_ids, + ClearContainersT = ClearContainersT{}) +{ + if constexpr(ClearContainersT::value) + { + callback_contexts.clear(); + buffered_contexts.clear(); + extern_corr_ids.clear(); + } + + const auto minimal_context_filter = [](const context_t* ctx) { + return (ctx->callback_tracer || ctx->buffered_tracer); + }; + + for(const auto* itr : context::get_active_contexts(minimal_context_filter)) + { + if(!itr) continue; + + // if the given domain + op is not enabled, skip this context + if(context_filter(itr, callback_domain_idx, operation_idx)) + { + callback_contexts.emplace_back( + callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); + extern_corr_ids.emplace(itr, empty_user_data); + } + + // if the given domain + op is not enabled, skip this context + if(context_filter(itr, buffered_domain_idx, operation_idx)) + { + buffered_contexts.emplace_back(buffered_context_data{itr}); + extern_corr_ids.emplace(itr, empty_user_data); + } + } +} + +template +inline void +populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, + rocprofiler_buffer_tracing_kind_t buffered_domain_idx, + callback_context_data_vec_t& callback_contexts, + buffered_context_data_vec_t& buffered_contexts, + external_correlation_id_map_t& extern_corr_ids, + ClearContainersT = ClearContainersT{}) +{ + if constexpr(ClearContainersT::value) + { + callback_contexts.clear(); + buffered_contexts.clear(); + extern_corr_ids.clear(); + } + + const auto minimal_context_filter = [](const context_t* ctx) { + return (ctx->callback_tracer || ctx->buffered_tracer); + }; + + for(const auto* itr : context::get_active_contexts(minimal_context_filter)) + { + if(!itr) continue; + + // if the given domain + op is not enabled, skip this context + if(context_filter(itr, callback_domain_idx)) + { + callback_contexts.emplace_back( + callback_context_data{itr, rocprofiler_callback_tracing_record_t{}}); + extern_corr_ids.emplace(itr, empty_user_data); + } + + // if the given domain + op is not enabled, skip this context + if(context_filter(itr, buffered_domain_idx)) + { + buffered_contexts.emplace_back(buffered_context_data{itr}); + extern_corr_ids.emplace(itr, empty_user_data); + } + } +} + +template +inline void +populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, + rocprofiler_buffer_tracing_kind_t buffered_domain_idx, + rocprofiler_tracing_operation_t operation_idx, + tracing_data& data, + ClearContainersT = ClearContainersT{}) +{ + populate_contexts(callback_domain_idx, + buffered_domain_idx, + operation_idx, + data.callback_contexts, + data.buffered_contexts, + data.external_correlation_ids); +} + +template +inline void +populate_contexts(rocprofiler_callback_tracing_kind_t callback_domain_idx, + rocprofiler_buffer_tracing_kind_t buffered_domain_idx, + tracing_data& data, + ClearContainersT = ClearContainersT{}) +{ + populate_contexts(callback_domain_idx, + buffered_domain_idx, + data.callback_contexts, + data.buffered_contexts, + data.external_correlation_ids); +} + +inline void +populate_external_correlation_ids(external_correlation_id_map_t& external_corr_ids, + rocprofiler_thread_id_t thr_id, + rocprofiler_external_correlation_id_request_kind_t kind, + rocprofiler_tracing_operation_t operation, + uint64_t internal_corr_id) +{ + for(auto& itr : external_corr_ids) + { + itr.second = itr.first->correlation_tracer.external_correlator.get( + thr_id, itr.first, kind, operation, internal_corr_id); + } +} + +inline void +update_external_correlation_ids(external_correlation_id_map_t& external_corr_ids, + rocprofiler_thread_id_t thr_id, + rocprofiler_external_correlation_id_request_kind_t kind) +{ + // enter callback may update the external correlation id field + for(auto& itr : external_corr_ids) + { + itr.second = + itr.first->correlation_tracer.external_correlator.update(itr.second, thr_id, kind); + } +} + +template +inline void +execute_phase_none_callbacks(callback_context_data_vec_t& callback_contexts, + rocprofiler_thread_id_t thr_id, + uint64_t internal_corr_id, + const external_correlation_id_map_t& external_corr_ids, + rocprofiler_callback_tracing_kind_t domain, + rocprofiler_tracing_operation_t operation, + TracerDataT& tracer_data) +{ + for(auto& itr : callback_contexts) + { + if(!context_filter(itr.ctx, domain, operation)) continue; + + auto& ctx = itr.ctx; + auto& record = itr.record; + auto& user_data = itr.user_data; + const auto& extern_corr_id_v = external_corr_ids.at(ctx); + + auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; + record = rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, + thr_id, + corr_id_v, + domain, + operation, + ROCPROFILER_CALLBACK_PHASE_NONE, + static_cast(&tracer_data)}; + + auto& callback_info = ctx->callback_tracer->callback_data.at(domain); + callback_info.callback(record, &user_data, callback_info.data); + } +} + +template +inline void +execute_phase_enter_callbacks(callback_context_data_vec_t& callback_contexts, + rocprofiler_thread_id_t thr_id, + uint64_t internal_corr_id, + const external_correlation_id_map_t& external_corr_ids, + rocprofiler_callback_tracing_kind_t domain, + rocprofiler_tracing_operation_t operation, + TracerDataT& tracer_data) +{ + for(auto& itr : callback_contexts) + { + if(!context_filter(itr.ctx, domain, operation)) continue; + + auto& ctx = itr.ctx; + auto& record = itr.record; + auto& user_data = itr.user_data; + const auto& extern_corr_id_v = external_corr_ids.at(ctx); + + auto corr_id_v = rocprofiler_correlation_id_t{internal_corr_id, extern_corr_id_v}; + record = rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, + thr_id, + corr_id_v, + domain, + operation, + ROCPROFILER_CALLBACK_PHASE_ENTER, + static_cast(&tracer_data)}; + + auto& callback_info = ctx->callback_tracer->callback_data.at(domain); + callback_info.callback(record, &user_data, callback_info.data); + } +} + +template +inline void +execute_phase_exit_callbacks(callback_context_data_vec_t& callback_contexts, + const external_correlation_id_map_t& external_corr_ids, + rocprofiler_callback_tracing_kind_t domain, + rocprofiler_tracing_operation_t operation, + TracerDataT& tracer_data) +{ + for(auto& itr : callback_contexts) + { + if(!context_filter(itr.ctx, domain, operation)) continue; + + auto& ctx = itr.ctx; + auto& record = itr.record; + auto& user_data = itr.user_data; + const auto& extern_corr_id_v = external_corr_ids.at(ctx); + + auto corr_id_v = + rocprofiler_correlation_id_t{record.correlation_id.internal, extern_corr_id_v}; + record = rocprofiler_callback_tracing_record_t{rocprofiler_context_id_t{ctx->context_idx}, + record.thread_id, + corr_id_v, + domain, + record.operation, + ROCPROFILER_CALLBACK_PHASE_EXIT, + static_cast(&tracer_data)}; + + auto& callback_info = ctx->callback_tracer->callback_data.at(domain); + callback_info.callback(record, &user_data, callback_info.data); + } +} + +template +inline void +execute_buffer_record_emplace(const buffered_context_data_vec_t& buffered_contexts, + rocprofiler_thread_id_t thr_id, + uint64_t internal_corr_id, + const external_correlation_id_map_t& external_corr_ids, + rocprofiler_buffer_tracing_kind_t domain, + OperationT operation, + BufferRecordT&& base_record) +{ + base_record.thread_id = thr_id; + base_record.kind = domain; + base_record.operation = operation; + // external correlation will be updated right before record is placed in buffer + base_record.correlation_id = rocprofiler_correlation_id_t{internal_corr_id, empty_user_data}; + + for(const auto& itr : buffered_contexts) + { + if(!context_filter(itr.ctx, domain, operation)) continue; + + auto buffer_id = itr.ctx->buffered_tracer->buffer_data.at(domain); + auto* buffer_v = buffer::get_buffer(buffer_id); + if(buffer_v && buffer_v->context_id == itr.ctx->context_idx && + buffer_v->buffer_id == buffer_id.handle) + { + // make copy of record + auto record_v = base_record; + // update the record with the correlation + record_v.correlation_id.external = external_corr_ids.at(itr.ctx); + + buffer_v->emplace(ROCPROFILER_BUFFER_CATEGORY_TRACING, domain, record_v); + } + } +} +} // namespace tracing +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py index 5069519548..7368ad7de2 100644 --- a/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py @@ -37,6 +37,7 @@ def test_data_structure(input_data): node_exists("hsa_api_traces", sdk_data["callback_records"]) node_exists("hip_api_traces", sdk_data["callback_records"], 0) node_exists("marker_api_traces", sdk_data["callback_records"]) + node_exists("kernel_dispatches", sdk_data["callback_records"]) node_exists("names", sdk_data["buffer_records"]) node_exists("kernel_dispatches", sdk_data["buffer_records"]) @@ -47,6 +48,32 @@ def test_data_structure(input_data): node_exists("retired_correlation_ids", sdk_data["buffer_records"]) +def test_size_entries(input_data): + # check that size fields are > 0 but account for function arguments + # which are named "size" + def check_size(data, bt): + if "size" in data.keys(): + if isinstance(data["size"], str) and bt.endswith('["args"]'): + pass + else: + assert data["size"] > 0, f"origin: {bt}" + + # recursively check the entire data structure + def iterate_data(data, bt): + if isinstance(data, (list, tuple)): + for i, itr in enumerate(data): + if isinstance(itr, dict): + check_size(itr, f"{bt}[{i}]") + iterate_data(itr, f"{bt}[{i}]") + elif isinstance(data, dict): + check_size(data, f"{bt}") + for key, itr in data.items(): + iterate_data(itr, f'{bt}["{key}"]') + + # start recursive check over entire JSON dict + iterate_data(input_data, "input_data") + + def test_timestamps(input_data): data = input_data sdk_data = data["rocprofiler-sdk-json-tool"] @@ -70,18 +97,26 @@ def test_timestamps(input_data): for titr in ["kernel_dispatches", "memory_copies"]: for itr in sdk_data["buffer_records"][titr]: - assert itr["start_timestamp"] < itr["end_timestamp"] - assert itr["correlation_id"]["internal"] > 0 - assert itr["correlation_id"]["external"] > 0 - assert sdk_data["metadata"]["init_time"] < itr["start_timestamp"] - assert sdk_data["metadata"]["init_time"] < itr["end_timestamp"] - assert sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] - assert sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}" + assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["end_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + ), f"[{titr}] {itr}" api_start = cb_start[itr["correlation_id"]["internal"]] api_end = cb_end[itr["correlation_id"]["internal"]] - assert api_start < itr["start_timestamp"] - assert api_end <= itr["end_timestamp"] + assert api_start < itr["start_timestamp"], f"[{titr}] {itr}" + assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" def test_internal_correlation_ids(input_data): @@ -126,18 +161,17 @@ def test_external_correlation_ids(input_data): extern_corr_ids = list(set(sorted(extern_corr_ids))) for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: for itr in sdk_data["buffer_records"][titr]: - assert itr["correlation_id"]["external"] > 0 - assert itr["thread_id"] == itr["correlation_id"]["external"] - assert itr["thread_id"] in extern_corr_ids - assert itr["correlation_id"]["external"] in extern_corr_ids + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + itr["thread_id"] == itr["correlation_id"]["external"] + ), f"[{titr}] {itr}" + assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" - for itr in sdk_data["buffer_records"]["kernel_dispatches"]: - assert itr["correlation_id"]["external"] > 0 - assert itr["correlation_id"]["external"] in extern_corr_ids - - for itr in sdk_data["buffer_records"]["memory_copies"]: - assert itr["correlation_id"]["external"] > 0 - assert itr["correlation_id"]["external"] in extern_corr_ids + for titr in ["kernel_dispatches", "memory_copies"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" def test_kernel_ids(input_data): @@ -162,6 +196,47 @@ def test_kernel_ids(input_data): for itr in sdk_data["buffer_records"]["kernel_dispatches"]: assert itr["kernel_id"] in symbol_info.keys() + for itr in sdk_data["callback_records"]["kernel_dispatches"]: + assert itr["payload"]["kernel_id"] in symbol_info.keys() + + +def test_kernel_dispatch_ids(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + num_dispatches = len(sdk_data["buffer_records"]["kernel_dispatches"]) + num_cb_dispatches = len(sdk_data["callback_records"]["kernel_dispatches"]) + + assert num_cb_dispatches == (2 * num_dispatches) + + bf_seq_ids = [] + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: + bf_seq_ids.append(itr["dispatch_id"]) + + cb_seq_ids = [] + for itr in sdk_data["callback_records"]["kernel_dispatches"]: + cb_seq_ids.append(itr["payload"]["dispatch_id"]) + + bf_seq_ids = sorted(bf_seq_ids) + cb_seq_ids = sorted(cb_seq_ids) + + assert (2 * len(bf_seq_ids)) == len(cb_seq_ids) + + assert bf_seq_ids[0] == cb_seq_ids[0] + assert bf_seq_ids[-1] == cb_seq_ids[-1] + + def get_uniq(data): + return list(set(data)) + + bf_seq_ids_uniq = get_uniq(bf_seq_ids) + cb_seq_ids_uniq = get_uniq(cb_seq_ids) + + assert bf_seq_ids == bf_seq_ids_uniq + assert len(cb_seq_ids) == (2 * len(cb_seq_ids_uniq)) + assert len(bf_seq_ids) == num_dispatches + assert len(bf_seq_ids_uniq) == num_dispatches + assert len(cb_seq_ids_uniq) == num_dispatches + def test_async_copy_direction(input_data): data = input_data @@ -223,13 +298,15 @@ def test_retired_correlation_ids(input_data): for cid, itr in async_corr_ids.items(): assert cid in retired_corr_ids.keys() - ts = retired_corr_ids[cid]["timestamp"] - assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}" + retired_ts = retired_corr_ids[cid]["timestamp"] + end_ts = itr["end_timestamp"] + assert (retired_ts - end_ts) > 0, f"correlation-id: {cid}, data: {itr}" for cid, itr in api_corr_ids.items(): assert cid in retired_corr_ids.keys() - ts = retired_corr_ids[cid]["timestamp"] - assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}" + retired_ts = retired_corr_ids[cid]["timestamp"] + end_ts = itr["end_timestamp"] + assert (retired_ts - end_ts) > 0, f"correlation-id: {cid}, data: {itr}" assert len(api_corr_ids.keys()) == (len(retired_corr_ids.keys())) diff --git a/projects/rocprofiler-sdk/tests/common/serialization.hpp b/projects/rocprofiler-sdk/tests/common/serialization.hpp index 882e584ff6..33d7177002 100644 --- a/projects/rocprofiler-sdk/tests/common/serialization.hpp +++ b/projects/rocprofiler-sdk/tests/common/serialization.hpp @@ -268,6 +268,39 @@ save(ArchiveT& ar, rocprofiler_callback_tracing_scratch_memory_data_t data) SAVE_DATA_FIELD(args_kind); } +template +void +save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data) +{ + SAVE_DATA_FIELD(size); + SAVE_DATA_FIELD(start_timestamp); + SAVE_DATA_FIELD(end_timestamp); + SAVE_DATA_FIELD(agent_id); + SAVE_DATA_FIELD(queue_id); + SAVE_DATA_FIELD(kernel_id); + SAVE_DATA_FIELD(dispatch_id); + SAVE_DATA_FIELD(private_segment_size); + SAVE_DATA_FIELD(group_segment_size); + SAVE_DATA_FIELD(workgroup_size); + SAVE_DATA_FIELD(group_segment_size); +} + +template +void +save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_data_t data) +{ + SAVE_DATA_FIELD(size); + SAVE_DATA_FIELD(agent_id); + SAVE_DATA_FIELD(queue_id); + SAVE_DATA_FIELD(kernel_id); + SAVE_DATA_FIELD(dispatch_id); + SAVE_DATA_FIELD(correlation_id); + SAVE_DATA_FIELD(private_segment_size); + SAVE_DATA_FIELD(group_segment_size); + SAVE_DATA_FIELD(workgroup_size); + SAVE_DATA_FIELD(grid_size); +} + template void save(ArchiveT& ar, rocprofiler_callback_tracing_record_t data) @@ -329,12 +362,15 @@ save(ArchiveT& ar, rocprofiler_buffer_tracing_kernel_dispatch_record_t data) { SAVE_DATA_FIELD(size); SAVE_DATA_FIELD(kind); + SAVE_DATA_FIELD(operation); + SAVE_DATA_FIELD(thread_id); SAVE_DATA_FIELD(correlation_id); SAVE_DATA_FIELD(start_timestamp); SAVE_DATA_FIELD(end_timestamp); SAVE_DATA_FIELD(agent_id); SAVE_DATA_FIELD(queue_id); SAVE_DATA_FIELD(kernel_id); + SAVE_DATA_FIELD(dispatch_id); SAVE_DATA_FIELD(private_segment_size); SAVE_DATA_FIELD(group_segment_size); SAVE_DATA_FIELD(workgroup_size); diff --git a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py index f5b985364a..02f20aa3b9 100644 --- a/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/kernel-tracing/validate.py @@ -37,6 +37,7 @@ def test_data_structure(input_data): node_exists("hsa_api_traces", sdk_data["callback_records"]) node_exists("hip_api_traces", sdk_data["callback_records"], 0) node_exists("marker_api_traces", sdk_data["callback_records"]) + node_exists("kernel_dispatches", sdk_data["callback_records"]) node_exists("names", sdk_data["buffer_records"]) node_exists("kernel_dispatches", sdk_data["buffer_records"]) @@ -47,6 +48,32 @@ def test_data_structure(input_data): node_exists("retired_correlation_ids", sdk_data["buffer_records"]) +def test_size_entries(input_data): + # check that size fields are > 0 but account for function arguments + # which are named "size" + def check_size(data, bt): + if "size" in data.keys(): + if isinstance(data["size"], str) and bt.endswith('["args"]'): + pass + else: + assert data["size"] > 0, f"origin: {bt}" + + # recursively check the entire data structure + def iterate_data(data, bt): + if isinstance(data, (list, tuple)): + for i, itr in enumerate(data): + if isinstance(itr, dict): + check_size(itr, f"{bt}[{i}]") + iterate_data(itr, f"{bt}[{i}]") + elif isinstance(data, dict): + check_size(data, f"{bt}") + for key, itr in data.items(): + iterate_data(itr, f'{bt}["{key}"]') + + # start recursive check over entire JSON dict + iterate_data(input_data, "input_data") + + def test_timestamps(input_data): data = input_data sdk_data = data["rocprofiler-sdk-json-tool"] @@ -70,18 +97,26 @@ def test_timestamps(input_data): for titr in ["kernel_dispatches", "memory_copies"]: for itr in sdk_data["buffer_records"][titr]: - assert itr["start_timestamp"] < itr["end_timestamp"] - assert itr["correlation_id"]["internal"] > 0 - assert itr["correlation_id"]["external"] > 0 - assert sdk_data["metadata"]["init_time"] < itr["start_timestamp"] - assert sdk_data["metadata"]["init_time"] < itr["end_timestamp"] - assert sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] - assert sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}" + assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["end_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + ), f"[{titr}] {itr}" api_start = cb_start[itr["correlation_id"]["internal"]] api_end = cb_end[itr["correlation_id"]["internal"]] - assert api_start < itr["start_timestamp"] - assert api_end <= itr["end_timestamp"] + assert api_start < itr["start_timestamp"], f"[{titr}] {itr}" + assert api_end <= itr["end_timestamp"], f"[{titr}] {itr}" def test_total_runtime(input_data): @@ -157,13 +192,15 @@ def test_retired_correlation_ids(input_data): for cid, itr in async_corr_ids.items(): assert cid in retired_corr_ids.keys() - ts = retired_corr_ids[cid]["timestamp"] - assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}" + retired_ts = retired_corr_ids[cid]["timestamp"] + end_ts = itr["end_timestamp"] + assert (retired_ts - end_ts) > 0, f"correlation-id: {cid}, data: {itr}" for cid, itr in api_corr_ids.items(): assert cid in retired_corr_ids.keys() - ts = retired_corr_ids[cid]["timestamp"] - assert (ts - itr["end_timestamp"]) > 0, f"correlation-id: {cid}, data: {itr}" + retired_ts = retired_corr_ids[cid]["timestamp"] + end_ts = itr["end_timestamp"] + assert (retired_ts - end_ts) > 0, f"correlation-id: {cid}, data: {itr}" assert len(api_corr_ids.keys()) == (len(retired_corr_ids.keys())) @@ -184,18 +221,17 @@ def test_external_correlation_ids(input_data): extern_corr_ids = list(set(sorted(extern_corr_ids))) for titr in ["hsa_api_traces", "marker_api_traces", "hip_api_traces"]: for itr in sdk_data["buffer_records"][titr]: - assert itr["correlation_id"]["external"] > 0 - assert itr["thread_id"] == itr["correlation_id"]["external"] - assert itr["thread_id"] in extern_corr_ids - assert itr["correlation_id"]["external"] in extern_corr_ids + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + itr["thread_id"] == itr["correlation_id"]["external"] + ), f"[{titr}] {itr}" + assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" - for itr in sdk_data["buffer_records"]["kernel_dispatches"]: - assert itr["correlation_id"]["external"] > 0 - assert itr["correlation_id"]["external"] in extern_corr_ids - - for itr in sdk_data["buffer_records"]["memory_copies"]: - assert itr["correlation_id"]["external"] > 0 - assert itr["correlation_id"]["external"] in extern_corr_ids + for titr in ["kernel_dispatches", "memory_copies"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" def test_kernel_ids(input_data): @@ -220,6 +256,47 @@ def test_kernel_ids(input_data): for itr in sdk_data["buffer_records"]["kernel_dispatches"]: assert itr["kernel_id"] in symbol_info.keys() + for itr in sdk_data["callback_records"]["kernel_dispatches"]: + assert itr["payload"]["kernel_id"] in symbol_info.keys() + + +def test_kernel_dispatch_ids(input_data): + data = input_data + sdk_data = data["rocprofiler-sdk-json-tool"] + + num_dispatches = len(sdk_data["buffer_records"]["kernel_dispatches"]) + num_cb_dispatches = len(sdk_data["callback_records"]["kernel_dispatches"]) + + assert num_cb_dispatches == (2 * num_dispatches) + + bf_seq_ids = [] + for itr in sdk_data["buffer_records"]["kernel_dispatches"]: + bf_seq_ids.append(itr["dispatch_id"]) + + cb_seq_ids = [] + for itr in sdk_data["callback_records"]["kernel_dispatches"]: + cb_seq_ids.append(itr["payload"]["dispatch_id"]) + + bf_seq_ids = sorted(bf_seq_ids) + cb_seq_ids = sorted(cb_seq_ids) + + assert (2 * len(bf_seq_ids)) == len(cb_seq_ids) + + assert bf_seq_ids[0] == cb_seq_ids[0] + assert bf_seq_ids[-1] == cb_seq_ids[-1] + + def get_uniq(data): + return list(set(data)) + + bf_seq_ids_uniq = get_uniq(bf_seq_ids) + cb_seq_ids_uniq = get_uniq(cb_seq_ids) + + assert bf_seq_ids == bf_seq_ids_uniq + assert len(cb_seq_ids) == (2 * len(cb_seq_ids_uniq)) + assert len(bf_seq_ids) == num_dispatches + assert len(bf_seq_ids_uniq) == num_dispatches + assert len(cb_seq_ids_uniq) == num_dispatches + def test_async_copy_direction(input_data): data = input_data diff --git a/projects/rocprofiler-sdk/tests/scratch-memory-tracing/validate.py b/projects/rocprofiler-sdk/tests/scratch-memory-tracing/validate.py index 238f28578f..0ffcaab10a 100755 --- a/projects/rocprofiler-sdk/tests/scratch-memory-tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/scratch-memory-tracing/validate.py @@ -76,13 +76,21 @@ def test_timestamps(input_data): for titr in ["kernel_dispatches", "memory_copies"]: for itr in sdk_data["buffer_records"][titr]: - assert itr["start_timestamp"] < itr["end_timestamp"] - assert itr["correlation_id"]["internal"] > 0 - assert itr["correlation_id"]["external"] > 0 - assert sdk_data["metadata"]["init_time"] < itr["start_timestamp"] - assert sdk_data["metadata"]["init_time"] < itr["end_timestamp"] - assert sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] - assert sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}" + assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["init_time"] < itr["end_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["start_timestamp"] + ), f"[{titr}] {itr}" + assert ( + sdk_data["metadata"]["fini_time"] > itr["end_timestamp"] + ), f"[{titr}] {itr}" # TODO(Is this check applicable for scratch, which doesn't use any correlation id?) # api_start = cb_start[itr["correlation_id"]["internal"]] @@ -133,22 +141,20 @@ def test_external_correlation_ids(input_data): extern_corr_ids = list(set(sorted(extern_corr_ids))) for titr in ["hsa_api_traces", "hip_api_traces"]: for itr in sdk_data["buffer_records"][titr]: - assert itr["correlation_id"]["external"] > 0 - assert itr["thread_id"] == itr["correlation_id"]["external"] - assert itr["thread_id"] in extern_corr_ids - assert itr["correlation_id"]["external"] in extern_corr_ids + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert ( + itr["thread_id"] == itr["correlation_id"]["external"] + ), f"[{titr}] {itr}" + assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" - for itr in sdk_data["buffer_records"]["kernel_dispatches"]: - assert itr["correlation_id"]["external"] > 0 - assert itr["correlation_id"]["external"] in extern_corr_ids - - for itr in sdk_data["buffer_records"]["memory_copies"]: - assert itr["correlation_id"]["external"] > 0 - assert itr["correlation_id"]["external"] in extern_corr_ids + for titr in ["kernel_dispatches", "memory_copies"]: + for itr in sdk_data["buffer_records"][titr]: + assert itr["correlation_id"]["external"] > 0, f"[{titr}] {itr}" + assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}" def op_name(op_name, record): - found_op = False op_key = None for kind_node in record["names"]["kind_names"]: @@ -159,6 +165,8 @@ def op_name(op_name, record): if op_node["key"] == op_key: return op_node + return None + # Tests above are identical to async-copy. Update as needed @@ -168,8 +176,10 @@ def test_scratch_memory_tracking(input_data): callback_records = sdk_data["callback_records"] buffer_records = sdk_data["buffer_records"] - scratch_callback_data = sdk_data["callback_records"]["scratch_memory_traces"] - scratch_buffer_data = sdk_data["buffer_records"]["scratch_memory_traces"] + scratch_callback_data = callback_records["scratch_memory_traces"] + scratch_buffer_data = buffer_records["scratch_memory_traces"] + + assert len(scratch_callback_data) == 2 * len(scratch_buffer_data) cb_op_names = op_name("SCRATCH_MEMORY", callback_records)["value"] bf_op_names = op_name("SCRATCH_MEMORY", buffer_records)["value"] @@ -226,14 +236,23 @@ def test_scratch_memory_tracking(input_data): for thread_id, nodes in cb_threads.items(): assert thread_id > 0 + # sort based on timestamp + nodes = sorted(nodes, key=lambda x: x["timestamp"]) + # start must be followed by end for inx in range(0, len(nodes), 2): this_node = nodes[inx] next_node = nodes[inx + 1] - assert rc(this_node)["phase"] + 1 == rc(next_node)["phase"] - assert rc(this_node)["thread_id"] == rc(next_node)["thread_id"] - assert this_node["timestamp"] < next_node["timestamp"] + assert ( + rc(this_node)["phase"] + 1 == rc(next_node)["phase"] + ), f"this:\n{this_node}\n\nnext:\n{next_node}" + assert ( + rc(this_node)["thread_id"] == rc(next_node)["thread_id"] + ), f"this:\n{this_node}\n\nnext:\n{next_node}" + assert ( + this_node["timestamp"] < next_node["timestamp"] + ), f"this:\n{this_node}\n\nnext:\n{next_node}" # alloc has more data vs free and async reclaim scratch_alloc_node = ( diff --git a/projects/rocprofiler-sdk/tests/tools/c-tool.c b/projects/rocprofiler-sdk/tests/tools/c-tool.c index c81c7a8894..dbe4cf2570 100644 --- a/projects/rocprofiler-sdk/tests/tools/c-tool.c +++ b/projects/rocprofiler-sdk/tests/tools/c-tool.c @@ -26,7 +26,7 @@ #endif /** - * @file samples/c_tool/client.cpp + * @file tests/tools/c-tool.c * * @brief Example rocprofiler client (tool) written in C */ diff --git a/projects/rocprofiler-sdk/tests/tools/json-tool.cpp b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp index 71dab5ac8a..bfdfb4cb9e 100644 --- a/projects/rocprofiler-sdk/tests/tools/json-tool.cpp +++ b/projects/rocprofiler-sdk/tests/tools/json-tool.cpp @@ -185,6 +185,14 @@ struct source_location } }; +template +auto +make_array(Tp&& arg, Args&&... args) +{ + constexpr auto N = sizeof...(Args) + 1; + return std::array{std::forward(arg), std::forward(args)...}; +} + using call_stack_t = std::vector; using buffer_kind_names_t = std::map; using buffer_kind_operation_names_t = @@ -490,6 +498,21 @@ struct marker_api_callback_record_t } }; +struct kernel_dispatch_callback_record_t +{ + uint64_t timestamp = 0; + rocprofiler_callback_tracing_record_t record = {}; + rocprofiler_callback_tracing_kernel_dispatch_data_t payload = {}; + + template + void save(ArchiveT& ar) const + { + ar(cereal::make_nvp("timestamp", timestamp)); + ar(cereal::make_nvp("record", record)); + ar(cereal::make_nvp("payload", payload)); + } +}; + struct scratch_memory_callback_record_t { uint64_t timestamp = 0; @@ -532,35 +555,21 @@ auto marker_api_cb_records = std::deque{}; auto counter_collection_bf_records = std::deque{}; auto hip_api_cb_records = std::deque{}; auto scratch_memory_cb_records = std::deque{}; +auto kernel_dispatch_cb_records = std::deque{}; -rocprofiler_thread_id_t -push_external_correlation(); - -void -counter_collection_buffered(rocprofiler_context_id_t, /*context*/ - rocprofiler_buffer_id_t, /*buffer_id*/ - rocprofiler_record_header_t** headers, - size_t num_headers, - void*, /*user_data*/ - uint64_t /*drop_count*/) +int +set_external_correlation_id(rocprofiler_thread_id_t thr_id, + rocprofiler_context_id_t ctx_id, + rocprofiler_external_correlation_id_request_kind_t kind, + rocprofiler_tracing_operation_t op, + uint64_t internal_corr_id, + rocprofiler_user_data_t* external_corr_id, + void* user_data) { - if(num_headers == 0) - throw std::runtime_error{"rocprofiler invoked a buffer callback with no headers " - "this should never happen"}; + consume_args(ctx_id, kind, op, internal_corr_id, user_data); - else if(headers == nullptr) - throw std::runtime_error{"rocprofiler invoked a buffer callback with a null pointer to the " - "array of headers. this should never happen"}; - - for(size_t i = 0; i < num_headers; ++i) - { - auto* header = headers[i]; - if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0) - { - auto* profiler_record = static_cast(header->payload); - counter_collection_bf_records.emplace_back(*profiler_record); - } - } + external_corr_id->value = thr_id; + return 0; } void @@ -653,13 +662,6 @@ tool_tracing_callback(rocprofiler_callback_tracing_record_t record, auto ts = rocprofiler_timestamp_t{}; ROCPROFILER_CALL(rocprofiler_get_timestamp(&ts), "get timestamp"); - static thread_local auto _once = std::once_flag{}; - std::call_once(_once, [&record]() { - // account for the fact that we are not wrapping pthread_create so the - // first external correlation id on a thread wont have updated value - record.correlation_id.external.value = push_external_correlation(); - }); - if(record.kind == ROCPROFILER_CALLBACK_TRACING_CODE_OBJECT) { if(record.operation == ROCPROFILER_CODE_OBJECT_LOAD) @@ -734,18 +736,28 @@ tool_tracing_callback(rocprofiler_callback_tracing_record_t record, auto _lk = std::unique_lock{_mutex}; scratch_memory_cb_records.emplace_back(scratch_memory_callback_record_t{ts, record, *data}); } + else if(record.kind == ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH) + { + auto* data = + static_cast(record.payload); + + static auto _mutex = std::mutex{}; + auto _lk = std::unique_lock{_mutex}; + kernel_dispatch_cb_records.emplace_back( + kernel_dispatch_callback_record_t{ts, record, *data}); + } else { throw std::runtime_error{"unsupported callback kind"}; } } -auto hsa_api_bf_records = std::deque{}; -auto marker_api_bf_records = std::deque{}; -auto hip_api_bf_records = std::deque{}; -auto kernel_dispatch_records = std::deque{}; -auto memory_copy_records = std::deque{}; -auto scratch_memory_records = std::deque{}; +auto hsa_api_bf_records = std::deque{}; +auto marker_api_bf_records = std::deque{}; +auto hip_api_bf_records = std::deque{}; +auto kernel_dispatch_bf_records = std::deque{}; +auto memory_copy_records = std::deque{}; +auto scratch_memory_records = std::deque{}; auto corr_id_retire_records = std::deque{}; @@ -757,13 +769,6 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/, void* user_data, uint64_t drop_count) { - // std::cerr << "[" << getpid() << "][" << __FUNCTION__ << "] buffer flush callback for " - // << num_headers << " records...\n" - // << std::flush; - - static auto _mutex = std::mutex{}; - auto _lk = std::unique_lock{_mutex}; - assert(user_data != nullptr); assert(drop_count == 0 && "drop count should be zero for lossless policy"); @@ -821,7 +826,7 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/, auto* record = static_cast( header->payload); - kernel_dispatch_records.emplace_back(*record); + kernel_dispatch_bf_records.emplace_back(*record); } else if(header->kind == ROCPROFILER_BUFFER_TRACING_MEMORY_COPY) { @@ -851,6 +856,11 @@ tool_tracing_buffered(rocprofiler_context_id_t /*context*/, "unexpected rocprofiler_record_header_t tracing category kind"}; } } + else if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS && header->kind == 0) + { + auto* profiler_record = static_cast(header->payload); + counter_collection_bf_records.emplace_back(*profiler_record); + } else { throw std::runtime_error{"unexpected rocprofiler_record_header_t category + kind"}; @@ -897,19 +907,26 @@ stop(); void flush(); +void +push_external_correlation(uint64_t _value); + +void +pop_external_correlation(); + // contexts -rocprofiler_context_id_t hsa_api_callback_ctx = {}; -rocprofiler_context_id_t hip_api_callback_ctx = {}; -rocprofiler_context_id_t marker_api_callback_ctx = {}; -rocprofiler_context_id_t code_object_ctx = {}; -rocprofiler_context_id_t hsa_api_buffered_ctx = {}; -rocprofiler_context_id_t hip_api_buffered_ctx = {}; -rocprofiler_context_id_t marker_api_buffered_ctx = {}; -rocprofiler_context_id_t kernel_dispatch_ctx = {}; -rocprofiler_context_id_t memory_copy_ctx = {}; -rocprofiler_context_id_t counter_collection_ctx = {}; -rocprofiler_context_id_t scratch_memory_ctx = {}; -rocprofiler_context_id_t corr_id_retire_ctx = {}; +rocprofiler_context_id_t hsa_api_callback_ctx = {}; +rocprofiler_context_id_t hip_api_callback_ctx = {}; +rocprofiler_context_id_t marker_api_callback_ctx = {}; +rocprofiler_context_id_t code_object_ctx = {}; +rocprofiler_context_id_t hsa_api_buffered_ctx = {}; +rocprofiler_context_id_t hip_api_buffered_ctx = {}; +rocprofiler_context_id_t marker_api_buffered_ctx = {}; +rocprofiler_context_id_t memory_copy_ctx = {}; +rocprofiler_context_id_t counter_collection_ctx = {}; +rocprofiler_context_id_t scratch_memory_ctx = {}; +rocprofiler_context_id_t corr_id_retire_ctx = {}; +rocprofiler_context_id_t kernel_dispatch_callback_ctx = {}; +rocprofiler_context_id_t kernel_dispatch_buffered_ctx = {}; // buffers rocprofiler_buffer_id_t hsa_api_buffered_buffer = {}; rocprofiler_buffer_id_t hip_api_buffered_buffer = {}; @@ -925,10 +942,11 @@ auto contexts = std::unordered_map{ {"HIP_API_CALLBACK", &hip_api_callback_ctx}, {"MARKER_API_CALLBACK", &marker_api_callback_ctx}, {"CODE_OBJECT", &code_object_ctx}, + {"KERNEL_DISPATCH_CALLBACK", &kernel_dispatch_callback_ctx}, {"HSA_API_BUFFERED", &hsa_api_buffered_ctx}, {"HIP_API_BUFFERED", &hip_api_buffered_ctx}, {"MARKER_API_BUFFERED", &marker_api_buffered_ctx}, - {"KERNEL_DISPATCH", &kernel_dispatch_ctx}, + {"KERNEL_DISPATCH_BUFFERED", &kernel_dispatch_buffered_ctx}, {"MEMORY_COPY", &memory_copy_ctx}, {"COUNTER_COLLECTION", &counter_collection_ctx}, {"SCRATCH_MEMORY", &scratch_memory_ctx}, @@ -986,6 +1004,9 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) for(auto itr : contexts) { ROCPROFILER_CALL(rocprofiler_create_context(itr.second), "context creation"); + ROCPROFILER_CALL(rocprofiler_configure_external_correlation_id_request_service( + *itr.second, nullptr, 0, set_external_correlation_id, nullptr), + "external correlation id request service configure"); } for(auto itr : {ROCPROFILER_CALLBACK_TRACING_HSA_CORE_API, @@ -1005,7 +1026,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) 0, tool_tracing_callback, nullptr), - "hip api callback tracing service configure"); + "hip runtime api callback tracing service configure"); ROCPROFILER_CALL( rocprofiler_configure_callback_tracing_service(code_object_ctx, @@ -1023,7 +1044,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) 0, tool_tracing_callback, nullptr), - "hsa api tracing service configure"); + "marker core api tracing service configure"); ROCPROFILER_CALL(rocprofiler_configure_callback_tracing_service( marker_api_callback_ctx, @@ -1032,7 +1053,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) 0, tool_tracing_callback, nullptr), - "hsa api tracing service configure"); + "marker control api tracing service configure"); ROCPROFILER_CALL( rocprofiler_configure_callback_tracing_service(marker_api_callback_ctx, @@ -1041,7 +1062,19 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) 0, tool_tracing_callback, nullptr), - "hsa api tracing service configure"); + "marker name api tracing service configure"); + + auto kernel_dispatch_cb_ops = + make_array(ROCPROFILER_KERNEL_DISPATCH_ENQUEUE); + + ROCPROFILER_CALL( + rocprofiler_configure_callback_tracing_service(kernel_dispatch_callback_ctx, + ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH, + kernel_dispatch_cb_ops.data(), + kernel_dispatch_cb_ops.size(), + tool_tracing_callback, + nullptr), + "kernel dispatch callback tracing service configure"); ROCPROFILER_CALL( rocprofiler_configure_callback_tracing_service(scratch_memory_ctx, @@ -1082,7 +1115,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) &marker_api_buffered_buffer), "buffer creation"); - ROCPROFILER_CALL(rocprofiler_create_buffer(kernel_dispatch_ctx, + ROCPROFILER_CALL(rocprofiler_create_buffer(kernel_dispatch_buffered_ctx, buffer_size, watermark, ROCPROFILER_BUFFER_POLICY_LOSSLESS, @@ -1122,8 +1155,8 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) buffer_size, watermark, ROCPROFILER_BUFFER_POLICY_LOSSLESS, - counter_collection_buffered, - nullptr, + tool_tracing_buffered, + tool_data, &counter_collection_buffer), "buffer creation"); @@ -1170,7 +1203,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) "buffer tracing service configure"); ROCPROFILER_CALL( - rocprofiler_configure_buffer_tracing_service(kernel_dispatch_ctx, + rocprofiler_configure_buffer_tracing_service(kernel_dispatch_buffered_ctx, ROCPROFILER_BUFFER_TRACING_KERNEL_DISPATCH, nullptr, 0, @@ -1249,9 +1282,15 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) options << "\n\t- " << itr.first; auto pos = context_settings.find(itr.first); if(pos == std::string::npos) + { + std::cerr << "Excluding context: " << itr.first << std::endl; itr.second = nullptr; + } else + { + std::cerr << "Enabling context: " << itr.first << std::endl; context_settings.erase(pos, itr.first.length()); + } } // detect if there are any invalid entries @@ -1261,7 +1300,8 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) auto msg = std::stringstream{}; msg << "[rocprofiler-sdk-json-tool][" << filename.substr(filename.find_last_of('/') + 1) << ":" << __LINE__ << "] invalid specification of ROCPROFILER_TOOL_CONTEXTS ('" - << context_settings_env << "'). Valid choices are: " << options.str(); + << context_settings_env << "'). Valid choices are: " << options.str() + << "\nRemainder: " << context_settings << "\n"; throw std::runtime_error{msg.str()}; } } @@ -1281,7 +1321,12 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) { options << "\n\t- " << itr.first; auto pos = context_settings.find(itr.first); - if(pos != std::string::npos) itr.second = nullptr; + if(pos != std::string::npos) + { + std::cerr << "Excluding context: " << itr.first << std::endl; + itr.second = nullptr; + context_settings.erase(pos, itr.first.length()); + } } // detect if there are any invalid entries @@ -1292,11 +1337,14 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) msg << "[rocprofiler-sdk-json-tool][" << filename.substr(filename.find_last_of('/') + 1) << ":" << __LINE__ << "] invalid specification of ROCPROFILER_TOOL_CONTEXTS_EXCLUDE ('" - << context_settings_excl_env << "'). Valid choices are: " << options.str(); + << context_settings_excl_env << "'). Valid choices are: " << options.str() + << "\nRemainder: " << context_settings << "\n"; throw std::runtime_error{msg.str()}; } } + push_external_correlation(getpid()); + start(); // no errors @@ -1318,6 +1366,8 @@ tool_fini(void* tool_data) stop(); flush(); + pop_external_correlation(); + rocprofiler_get_timestamp(&fini_time); std::cerr << "[" << getpid() << "][" << __FUNCTION__ @@ -1327,14 +1377,16 @@ tool_fini(void* tool_data) << ", hsa_api_callback_records=" << hsa_api_cb_records.size() << ", hip_api_callback_records=" << hip_api_cb_records.size() << ", marker_api_callback_records=" << marker_api_cb_records.size() - << ", kernel_dispatch_records=" << kernel_dispatch_records.size() + << ", scratch_memory_callback_records=" << scratch_memory_cb_records.size() + << ", kernel_dispatch_callback_records=" << kernel_dispatch_cb_records.size() + << ", kernel_dispatch_bf_records=" << kernel_dispatch_bf_records.size() << ", memory_copy_records=" << memory_copy_records.size() << ", scratch_memory_records=" << scratch_memory_records.size() << ", hsa_api_bf_records=" << hsa_api_bf_records.size() << ", hip_api_bf_records=" << hip_api_bf_records.size() << ", marker_api_bf_records=" << marker_api_bf_records.size() << ", corr_id_retire_records=" << corr_id_retire_records.size() - << ", counter_collection_records" << counter_collection_bf_records.size() << "...\n" + << ", counter_collection_records=" << counter_collection_bf_records.size() << "...\n" << std::flush; auto* _call_stack = static_cast(tool_data); @@ -1418,6 +1470,7 @@ write_json(call_stack_t* _call_stack) json_ar(cereal::make_nvp("hip_api_traces", hip_api_cb_records)); json_ar(cereal::make_nvp("marker_api_traces", marker_api_cb_records)); json_ar(cereal::make_nvp("scratch_memory_traces", scratch_memory_cb_records)); + json_ar(cereal::make_nvp("kernel_dispatches", kernel_dispatch_cb_records)); } catch(std::exception& e) { std::cerr << "[" << getpid() << "][" << __FUNCTION__ @@ -1431,7 +1484,7 @@ write_json(call_stack_t* _call_stack) try { json_ar(cereal::make_nvp("names", buffer_name_info)); - json_ar(cereal::make_nvp("kernel_dispatches", kernel_dispatch_records)); + json_ar(cereal::make_nvp("kernel_dispatches", kernel_dispatch_bf_records)); json_ar(cereal::make_nvp("memory_copies", memory_copy_records)); json_ar(cereal::make_nvp("scratch_memory_traces", scratch_memory_records)); json_ar(cereal::make_nvp("hsa_api_traces", hsa_api_bf_records)); @@ -1512,7 +1565,7 @@ write_perfetto() agent_ids.emplace(itr.src_agent_id.handle); } - for(auto itr : kernel_dispatch_records) + for(auto itr : kernel_dispatch_bf_records) agent_queue_ids[itr.agent_id.handle].emplace(itr.queue_id.handle); } @@ -1624,7 +1677,7 @@ write_perfetto() itr.kind, "operation", itr.operation, - "cid", + "corr_id", itr.correlation_id.internal, [&](::perfetto::EventContext ctx) { for(const auto& aitr : _args) @@ -1663,7 +1716,7 @@ write_perfetto() itr.kind, "operation", itr.operation, - "cid", + "corr_id", itr.correlation_id.internal, [&](::perfetto::EventContext ctx) { for(const auto& aitr : _args) @@ -1704,7 +1757,7 @@ write_perfetto() } auto demangled = std::unordered_map{}; - for(auto itr : kernel_dispatch_records) + for(auto itr : kernel_dispatch_bf_records) { const kernel_symbol_callback_record_t* sym = nullptr; for(const auto& kitr : kernel_symbol_records) @@ -1736,11 +1789,11 @@ write_perfetto() itr.kind, "agent", itr.agent_id.handle, - "cid", + "corr_id", itr.correlation_id.internal, "queue", itr.queue_id.handle, - "kid", + "kernel_id", itr.kernel_id, "private_segment_size", itr.private_segment_size, @@ -1827,8 +1880,8 @@ flush() } } -rocprofiler_thread_id_t -push_external_correlation() +void +push_external_correlation(uint64_t _value) { auto tid = rocprofiler_thread_id_t{}; ROCPROFILER_CALL(rocprofiler_get_thread_id(&tid), "get thread id"); @@ -1838,11 +1891,36 @@ push_external_correlation() if(itr.second) { ROCPROFILER_CALL(rocprofiler_push_external_correlation_id( - *itr.second, tid, rocprofiler_user_data_t{.value = tid}), + *itr.second, tid, rocprofiler_user_data_t{.value = _value}), "push external correlation"); } } - return tid; +} + +void +pop_external_correlation() +{ + auto tid = rocprofiler_thread_id_t{}; + ROCPROFILER_CALL(rocprofiler_get_thread_id(&tid), "get thread id"); + + for(auto itr : contexts) + { + if(itr.second) + { + auto _data = rocprofiler_user_data_t{.value = 0}; + ROCPROFILER_CALL(rocprofiler_pop_external_correlation_id(*itr.second, tid, &_data), + "push external correlation"); + if(_data.value != static_cast(getpid())) + { + auto _msg = std::stringstream{}; + _msg << "rocprofiler_pop_external_correlation_id(context.handle=" + << itr.second->handle << ", tid=" << tid + << ", ...) returned external correlation id value of " << _data.value + << ". expected: " << getpid(); + throw std::runtime_error{_msg.str()}; + } + } + } } } // namespace } // namespace client