Callback tracing for kernel dispatches + External correlation ID request service (#682)

* Support ROCPROFILER_CALLBACK_TRACING_KERNEL_DISPATCH

* Fix doxygen

* Update callback tracing

- temporary hacks for kind operation name and iterate kind operations

* Update source/include/rocprofiler-sdk

- introduce sequence id for kernel dispatches

* Update lib/rocprofiler-sdk (seq id)

- support sequence id passing

* Update tests (seq id)

- testing for sequence ids

* Cleanup include/rocprofiler-sdk/fwd.h

* Misc cleanup

* External Correlation ID Request Service (#699)

* External correlation ID request service

- callback requesting an external correlation ID instead of fetching from top of pushed external correlation ID stack

* Update external correlation id request support

- pass internal correlation ID in callback
- async copy generates a correlation ID if none already exists
- added external correlation ID request support for scratch memory tracing
- updated scratch memory tracing to use tracing:: functions

* Update hsa/queue.hpp

- new line at EOF

* Misc tweaks

- remove unnecessary logging in agent.cpp
- correlation_id::add_ref_count check for retirement
- finalization check in HSA queue AsyncSignalHandler

* Improve assertion failure logging in misc tests

* Update include/rocprofiler-sdk/fwd.h

- remove rocprofiler_record_counter_header_t

* Move lib/rocprofiler-sdk/tracing.hpp into lib/rocprofiler-sdk/tracing/ folder

* Update lib/rocprofiler-sdk/hsa/*

- hsa::get_hsa_status_string
- queue_info_session.hpp header
- rocprofiler_packet.hpp

* Update lib/rocprofiler-sdk/{counters,hip,marker}

- execute_phase_exit_callbacks tweaks
- queue_info_session tweaks

* Move rocprofiler_kernel_dispatch_operation_t to include/rocprofiler-sdk/fwd.h

* Update rocprofiler_buffer_tracing_kernel_dispatch_record_t

- add operation field and thread_id field

* Add lib/rocprofiler-sdk/kernel_dispatch

- enum <-> string mapping for kernel dispatch
- tracing implementations

* Update lib/rocprofiler-sdk/CMakeLists.txt

- tracing and kernel dispatch sub-directories

* Update lib/rocprofiler-sdk/{buffer,callback}_tracing.cpp

- invoke rocprofiler::kernel_tracing functions

* Update tests/common/serialization.hpp

- support operation and thread_id fields for rocprofiler_buffer_tracing_kernel_dispatch_record_t

* Update tests/tools/json-tool.cpp

- use external correlation id request service

* Rename sequence_id to dispatch_id

[ROCm/rocprofiler-sdk commit: 56030018dc]

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;
 
 /**

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;

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;
 
 /**

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;
 
 /** @} */
 

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;

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

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/agent.cpp

@@ -52,22 +52,14 @@
 #include <unordered_map>
 #include <vector>
 
-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<std::pair<size_t, std::unique_ptr<std::string>>>;
 
 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<hsa_agent_t>
 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;
     }
 

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:
         {

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 <rocprofiler-sdk/callback_tracing.h>
 #include <rocprofiler-sdk/fwd.h>
 #include <rocprofiler-sdk/hip/table_id.h>
 #include <rocprofiler-sdk/hsa/table_id.h>
@@ -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_kind_t>{
-        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;
         }
     };
 

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;

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.cpp

@@ -266,7 +266,8 @@ queue_cb(const context::context*                                         ctx,
          const std::shared_ptr<counter_callback_info>&                   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<counter_callback_info>& 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);

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/counters/core.hpp

@@ -130,7 +130,8 @@ queue_cb(const context::context*                                         ctx,
          const std::shared_ptr<counter_callback_info>&                   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);

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());

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<request_kind_t>& 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<rocprofiler_user_data_t>
+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<rocprofiler_external_correlation_id_request_kind_t>{};
+    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,

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 <bitset>
 #include <cstdint>
+#include <optional>
 #include <string>
 #include <unordered_map>
 #include <vector>
 
 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<request_kind_t>& kinds_v);
+
+    bool requires_request(request_kind_t kind) const;
+
 private:
-    common::Synchronized<external_correlation_map_t> data = {};
+    rocprofiler_user_data_t get(rocprofiler_thread_id_t thr_id) const;
+
+    std::optional<rocprofiler_user_data_t> 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_kind_size>                   request       = 0;
+    common::Synchronized<external_correlation_map_t> data          = {};
 };
 }  // namespace external_correlation
 }  // namespace rocprofiler

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 <rocprofiler-sdk/buffer.h>
 #include <rocprofiler-sdk/callback_tracing.h>
@@ -179,57 +180,6 @@ hip_api_impl<TableIdx, OpIdx>::exec(FuncT&& _func, Args&&... args)
     return get_default_retval<return_type>();
 }
 
-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_context_data>& callback_contexts,
-                  std::vector<buffered_context_data>& 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 <size_t TableIdx, size_t OpIdx>
 template <typename RetT, typename... Args>
 RetT
@@ -239,17 +189,23 @@ hip_api_impl<TableIdx, OpIdx>::functor(Args... args)
     using callback_api_data_t = typename hip_domain_info<TableIdx>::callback_data_type;
     using buffered_api_data_t = typename hip_domain_info<TableIdx>::buffered_data_type;
 
+    constexpr auto external_corr_id_domain_idx =
+        hip_domain_info<TableIdx>::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<callback_context_data>{};
-    auto buffered_contexts = std::vector<buffered_context_data>{};
+    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<TableIdx, OpIdx>::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<TableIdx, OpIdx>::functor(Args... args)
         set_data_args(info_type::get_api_data_args(tracer_data.args),
                       convert_arg_type(std::forward<Args>(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<void*>(&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<TableIdx, OpIdx>::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<void*>(&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

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 <rocprofiler-sdk/external_correlation.h>
 #    include <rocprofiler-sdk/hip.h>
 
 #    include "defines.hpp"
@@ -49,6 +50,8 @@ struct hip_domain_info<ROCPROFILER_HIP_TABLE_ID_Runtime>
     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<ROCPROFILER_HIP_TABLE_ID_Compiler>
     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

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)

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;

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 <rocprofiler-sdk/buffer.h>
 #include <rocprofiler-sdk/callback_tracing.h>
@@ -258,65 +259,17 @@ hsa_api_impl<TableIdx, OpIdx>::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_context_data>& callback_contexts,
-                  std::vector<buffered_context_data>& 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 <size_t TableIdx, size_t OpIdx>
 template <typename RetT, typename... Args>
 RetT
 hsa_api_impl<TableIdx, OpIdx>::functor(Args... args)
 {
-    using info_type = hsa_api_info<TableIdx, OpIdx>;
+    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<TableIdx, OpIdx>;
+
+    constexpr auto external_corr_id_domain_idx =
+        hsa_domain_info<TableIdx>::external_correlation_id_domain_idx;
 
     LOG_IF(INFO, registration::get_fini_status() != 0) << "Executing " << info_type::name;
 
@@ -329,15 +282,18 @@ hsa_api_impl<TableIdx, OpIdx>::functor(Args... args)
             return;
     }
 
-    auto thr_id            = common::get_tid();
-    auto callback_contexts = std::vector<callback_context_data>{};
-    auto buffered_contexts = std::vector<buffered_context_data>{};
+    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<TableIdx, OpIdx>::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>(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>(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<void*>(&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<TableIdx, OpIdx>::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<void*>(&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 <size_t TableIdx>
 const char*

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 <rocprofiler-sdk/external_correlation.h>
 #include <rocprofiler-sdk/fwd.h>
 #include <rocprofiler-sdk/hsa.h>
 #include <rocprofiler-sdk/hsa/table_id.h>
@@ -58,6 +59,8 @@ struct hsa_domain_info<ROCPROFILER_HSA_TABLE_ID_Core>
     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<ROCPROFILER_HSA_TABLE_ID_AmdExt>
     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<ROCPROFILER_HSA_TABLE_ID_ImageExt>
     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<ROCPROFILER_HSA_TABLE_ID_FinalizeExt>
     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<ROCPROFILER_HSA_TABLE_ID_AmdTool>
     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

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/hsa.hpp

@@ -26,6 +26,7 @@
 #include <rocprofiler-sdk/rocprofiler.h>
 
 #include <cstdint>
+#include <string_view>
 #include <vector>
 
 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 <size_t TableIdx>
 const char*
 name_by_id(uint32_t id);

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 <rocprofiler-sdk/callback_tracing.h>
+#include <rocprofiler-sdk/external_correlation.h>
+#include <rocprofiler-sdk/fwd.h>
 
 #include <glog/logging.h>
 #include <hsa/hsa.h>
 #include <hsa/hsa_ext_amd.h>
-#include <rocprofiler-sdk/fwd.h>
 
 #include <atomic>
 #include <chrono>
@@ -69,93 +76,49 @@ namespace hsa
 {
 namespace
 {
+template <typename DomainT, typename... Args>
+inline bool
+context_filter(const context::context* ctx, DomainT domain, Args... args)
+{
+    if constexpr(std::is_same<DomainT, rocprofiler_buffer_tracing_kind_t>::value)
+    {
+        return (ctx->buffered_tracer && ctx->buffered_tracer->domains(domain, args...));
+    }
+    else if constexpr(std::is_same<DomainT, rocprofiler_callback_tracing_kind_t>::value)
+    {
+        return (ctx->callback_tracer && ctx->callback_tracer->domains(domain, args...));
+    }
+    else
+    {
+        static_assert(common::mpl::assert_false<DomainT>::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<Queue::queue_info_session_t**>(data);
+        delete _session;
+        return false;
+    }
+
     auto& queue_info_session = *static_cast<Queue::queue_info_session_t*>(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<rocprofiler_dispatch_id_t>{0};
+
     auto&& CreateBarrierPacket = [](hsa_signal_t*                    dependency_signal,
                                     hsa_signal_t*                    completion_signal,
                                     std::vector<rocprofiler_packet>& _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<Queue*>(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<const rocprofiler_packet*>(packets);
     auto        transformed_packets = std::vector<rocprofiler_packet>{};
@@ -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

projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/queue.hpp

@@ -24,6 +24,7 @@
 
 #include <rocprofiler-sdk/agent.h>
 #include <rocprofiler-sdk/buffer_tracing.h>
+#include <rocprofiler-sdk/callback_tracing.h>
 #include <rocprofiler-sdk/fwd.h>
 
 #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 <hsa/amd_hsa_kernel_code.h>
 #include <hsa/hsa.h>
@@ -40,6 +43,7 @@
 #include <hsa/hsa_ven_amd_loader.h>
 
 #include <atomic>
+#include <cstdint>
 #include <functional>
 #include <iostream>
 #include <memory>
@@ -56,44 +60,6 @@ struct correlation_id;
 }  // namespace context
 namespace hsa
 {
-using ClientID = int64_t;
-using inst_pkt_t =
-    common::container::small_vector<std::pair<std::unique_ptr<AQLPacket>, 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<const context_t*>;
-    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<const context_t*, rocprofiler_user_data_t>;
-
-        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<const context_t*>;
+    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<std::unique_ptr<AQLPacket>(
         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<std::mutex> lock(_lock_queue);
     func();
 }
-
 }  // namespace hsa
-}  // namespace rocprofiler
+}  // namespace rocprofiler

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<rocprofiler::ThreadTracer>),

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 <rocprofiler-sdk/callback_tracing.h>
+#include <rocprofiler-sdk/fwd.h>
+
+#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<const context_t*, rocprofiler_user_data_t>;
+    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<const context_t*>;
+
+    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

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 <rocprofiler-sdk/callback_tracing.h>
+#include <rocprofiler-sdk/fwd.h>
+
+#include "lib/rocprofiler-sdk/hsa/aql_packet.hpp"
+
+#include <hsa/hsa_ven_amd_aqlprofile.h>
+#include <hsa/hsa_ven_amd_loader.h>
+
+namespace rocprofiler
+{
+namespace hsa
+{
+using ClientID = int64_t;
+
+using inst_pkt_t =
+    common::container::small_vector<std::pair<std::unique_ptr<AQLPacket>, 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

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 <rocprofiler-sdk/agent.h>
+#include <rocprofiler-sdk/external_correlation.h>
 #include <rocprofiler-sdk/fwd.h>
 #include <rocprofiler-sdk/hsa/api_id.h>
 #include <rocprofiler-sdk/hsa/table_id.h>
@@ -214,6 +216,8 @@ struct scratch_op_info<hsa_amd_tool_id_none>
         using end_fn_t   = scratch_op_info<hsa_amd_tool_id_##ENDPHASE>::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_context_data>& callback_contexts,
-                  std::vector<buffered_context_data>& 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_context_data> callback_ctxs = {};
-        std::vector<buffered_context_data> 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<ScratchOpIdx>::operation;
     constexpr auto OpPhase = scratch_op_info<ScratchOpIdx>::phase;
+    constexpr auto external_corr_id_domain_idx =
+        amd_tool_api_info<OpIdx>::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<OpIdx>(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<void*>(&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));
         }
     }
 

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})

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 <string_view>
+
+#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<ROCPROFILER_KERNEL_DISPATCH_##CODE>                                \
+    {                                                                                              \
+        static constexpr auto operation_idx = ROCPROFILER_KERNEL_DISPATCH_##CODE;                  \
+        static constexpr auto name          = #CODE;                                               \
+    };
+
+template <size_t Idx>
+struct kernel_dispatch_info;
+
+ROCPROFILER_KERNEL_DISPATCH_INFO(NONE)
+ROCPROFILER_KERNEL_DISPATCH_INFO(ENQUEUE)
+ROCPROFILER_KERNEL_DISPATCH_INFO(COMPLETE)
+
+template <size_t Idx, size_t... IdxTail>
+const char*
+name_by_id(const uint32_t id, std::index_sequence<Idx, IdxTail...>)
+{
+    if(Idx == id) return kernel_dispatch_info<Idx>::name;
+    if constexpr(sizeof...(IdxTail) > 0)
+        return name_by_id(id, std::index_sequence<IdxTail...>{});
+    else
+        return nullptr;
+}
+
+template <size_t Idx, size_t... IdxTail>
+uint32_t
+id_by_name(const char* name, std::index_sequence<Idx, IdxTail...>)
+{
+    if(std::string_view{kernel_dispatch_info<Idx>::name} == std::string_view{name})
+        return kernel_dispatch_info<Idx>::operation_idx;
+    if constexpr(sizeof...(IdxTail) > 0)
+        return id_by_name(name, std::index_sequence<IdxTail...>{});
+    else
+        return ROCPROFILER_HSA_AMD_EXT_API_ID_NONE;
+}
+
+template <size_t... Idx>
+void
+get_ids(std::vector<uint32_t>& _id_list, std::index_sequence<Idx...>)
+{
+    auto _emplace = [](auto& _vec, uint32_t _v) {
+        if(_v < static_cast<uint32_t>(ROCPROFILER_HSA_AMD_EXT_API_ID_LAST)) _vec.emplace_back(_v);
+    };
+
+    (_emplace(_id_list, kernel_dispatch_info<Idx>::operation_idx), ...);
+}
+
+template <size_t... Idx>
+void
+get_names(std::vector<const char*>& _name_list, std::index_sequence<Idx...>)
+{
+    auto _emplace = [](auto& _vec, const char* _v) {
+        if(_v != nullptr && strnlen(_v, 1) > 0) _vec.emplace_back(_v);
+    };
+
+    (_emplace(_name_list, kernel_dispatch_info<Idx>::name), ...);
+}
+}  // namespace
+
+const char*
+name_by_id(uint32_t id)
+{
+    return name_by_id(id, std::make_index_sequence<ROCPROFILER_KERNEL_DISPATCH_LAST>{});
+}
+
+uint32_t
+id_by_name(const char* name)
+{
+    return id_by_name(name, std::make_index_sequence<ROCPROFILER_KERNEL_DISPATCH_LAST>{});
+}
+
+std::vector<uint32_t>
+get_ids()
+{
+    auto _data = std::vector<uint32_t>{};
+    _data.reserve(ROCPROFILER_KERNEL_DISPATCH_LAST);
+    get_ids(_data, std::make_index_sequence<ROCPROFILER_KERNEL_DISPATCH_LAST>{});
+    return _data;
+}
+
+std::vector<const char*>
+get_names()
+{
+    auto _data = std::vector<const char*>{};
+    _data.reserve(ROCPROFILER_KERNEL_DISPATCH_LAST);
+    get_names(_data, std::make_index_sequence<ROCPROFILER_KERNEL_DISPATCH_LAST>{});
+    return _data;
+}
+}  // namespace kernel_dispatch
+}  // namespace rocprofiler

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 <rocprofiler-sdk/rocprofiler.h>
+
+#include <cstdint>
+#include <vector>
+
+namespace rocprofiler
+{
+namespace kernel_dispatch
+{
+const char*
+name_by_id(uint32_t id);
+
+uint32_t
+id_by_name(const char* name);
+
+std::vector<const char*>
+get_names();
+
+std::vector<uint32_t>
+get_ids();
+}  // namespace kernel_dispatch
+}  // namespace rocprofiler

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 <rocprofiler-sdk/callback_tracing.h>
+#include <rocprofiler-sdk/fwd.h>
+
+#include <hsa/hsa.h>
+
+#include <string_view>
+
+#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

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<const context_t*, rocprofiler_user_data_t>;
+using external_corr_id_map_t = user_data_map_t;
+
+void
+dispatch_complete(hsa::queue_info_session&);
+}  // namespace kernel_dispatch
+}  // namespace rocprofiler

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 <rocprofiler-sdk/buffer.h>
 #include <rocprofiler-sdk/callback_tracing.h>
@@ -130,57 +131,6 @@ roctx_api_impl<TableIdx, OpIdx>::exec(FuncT&& _func, Args&&... args)
         return get_default_retval<return_type>();
 }
 
-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_context_data>& callback_contexts,
-                  std::vector<buffered_context_data>& 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 <size_t TableIdx, size_t OpIdx>
 template <typename RetT, typename... Args>
 RetT
@@ -190,17 +140,22 @@ roctx_api_impl<TableIdx, OpIdx>::functor(Args... args)
     using callback_api_data_t = typename roctx_domain_info<TableIdx>::callback_data_type;
     using buffered_api_data_t = typename roctx_domain_info<TableIdx>::buffer_data_type;
 
+    constexpr auto external_corr_id_domain_idx =
+        roctx_domain_info<TableIdx>::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<callback_context_data>{};
-    auto buffered_contexts = std::vector<buffered_context_data>{};
+    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<TableIdx, OpIdx>::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>(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>(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<void*>(&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<TableIdx, OpIdx>::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<void*>(&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

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 <rocprofiler-sdk/external_correlation.h>
 #include <rocprofiler-sdk/marker/table_id.h>
 
 namespace rocprofiler
@@ -42,33 +43,39 @@ template <>
 struct roctx_domain_info<ROCPROFILER_MARKER_TABLE_ID_RoctxCore>
 : roctx_domain_info<ROCPROFILER_MARKER_TABLE_ID_LAST>
 {
+    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<ROCPROFILER_MARKER_TABLE_ID_RoctxControl>
 : roctx_domain_info<ROCPROFILER_MARKER_TABLE_ID_LAST>
 {
+    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<ROCPROFILER_MARKER_TABLE_ID_RoctxName>
 : roctx_domain_info<ROCPROFILER_MARKER_TABLE_ID_LAST>
 {
+    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

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}")

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,

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})

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 <rocprofiler-sdk/fwd.h>
+
+#include <functional>
+#include <unordered_map>
+#include <vector>
+
+namespace rocprofiler
+{
+namespace context
+{
+struct context;
+struct correlation_tracing_service;
+}  // namespace context
+namespace tracing
+{
+template <typename Tp, size_t N>
+using small_vector_t                = common::container::small_vector<Tp, N>;
+using correlation_service           = context::correlation_tracing_service;
+using context_t                     = context::context;
+using context_array_t               = common::container::small_vector<const context_t*>;
+using external_correlation_id_map_t = std::unordered_map<const context_t*, rocprofiler_user_data_t>;
+
+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<callback_context_data, context_data_vec_size>;
+using buffered_context_data_vec_t = small_vector_t<buffered_context_data, context_data_vec_size>;
+
+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

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 <rocprofiler-sdk/fwd.h>
+
+#include <functional>
+#include <unordered_map>
+#include <vector>
+
+namespace rocprofiler
+{
+namespace tracing
+{
+// template <typename DomainT, typename... Args>
+// bool
+// context_filter(const context::context* ctx, DomainT domain, Args... args);
+
+// template <typename ClearContainersT = std::false_type>
+// 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 <typename ClearContainersT = std::false_type>
+// 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 <typename ClearContainersT = std::false_type>
+// 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 <typename ClearContainersT = std::false_type>
+// 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 <typename TracerDataT>
+// 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 <typename TracerDataT>
+// 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 <typename TracerDataT>
+// 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 <typename BufferRecordT, typename OperationT = rocprofiler_tracing_operation_t>
+// 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 <typename DomainT, typename... Args>
+inline bool
+context_filter(const context::context* ctx, DomainT domain, Args... args)
+{
+    if constexpr(std::is_same<DomainT, rocprofiler_buffer_tracing_kind_t>::value)
+    {
+        return (ctx->buffered_tracer && ctx->buffered_tracer->domains(domain, args...));
+    }
+    else if constexpr(std::is_same<DomainT, rocprofiler_callback_tracing_kind_t>::value)
+    {
+        return (ctx->callback_tracer && ctx->callback_tracer->domains(domain, args...));
+    }
+    else
+    {
+        static_assert(common::mpl::assert_false<DomainT>::value, "unsupported domain type");
+        return false;
+    }
+}
+
+template <typename ClearContainersT = std::false_type>
+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 <typename ClearContainersT = std::false_type>
+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 <typename ClearContainersT = std::false_type>
+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<ClearContainersT>(callback_domain_idx,
+                                        buffered_domain_idx,
+                                        operation_idx,
+                                        data.callback_contexts,
+                                        data.buffered_contexts,
+                                        data.external_correlation_ids);
+}
+
+template <typename ClearContainersT = std::false_type>
+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<ClearContainersT>(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 <typename TracerDataT>
+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<void*>(&tracer_data)};
+
+        auto& callback_info = ctx->callback_tracer->callback_data.at(domain);
+        callback_info.callback(record, &user_data, callback_info.data);
+    }
+}
+
+template <typename TracerDataT>
+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<void*>(&tracer_data)};
+
+        auto& callback_info = ctx->callback_tracer->callback_data.at(domain);
+        callback_info.callback(record, &user_data, callback_info.data);
+    }
+}
+
+template <typename TracerDataT>
+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<void*>(&tracer_data)};
+
+        auto& callback_info = ctx->callback_tracer->callback_data.at(domain);
+        callback_info.callback(record, &user_data, callback_info.data);
+    }
+}
+
+template <typename BufferRecordT, typename OperationT = rocprofiler_tracing_operation_t>
+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

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()))
 

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 <typename ArchiveT>
+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 <typename ArchiveT>
+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 <typename ArchiveT>
 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);

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

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 = (

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
  */

projects/rocprofiler-sdk/tests/tools/json-tool.cpp

@@ -185,6 +185,14 @@ struct source_location
     }
 };
 
+template <typename Tp, typename... Args>
+auto
+make_array(Tp&& arg, Args&&... args)
+{
+    constexpr auto N = sizeof...(Args) + 1;
+    return std::array<Tp, N>{std::forward<Tp>(arg), std::forward<Args>(args)...};
+}
+
 using call_stack_t        = std::vector<source_location>;
 using buffer_kind_names_t = std::map<rocprofiler_buffer_tracing_kind_t, std::string>;
 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 <typename ArchiveT>
+    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<marker_api_callback_record_t>{};
 auto counter_collection_bf_records = std::deque<rocprofiler_record_counter_t>{};
 auto hip_api_cb_records            = std::deque<hip_api_callback_record_t>{};
 auto scratch_memory_cb_records     = std::deque<scratch_memory_callback_record_t>{};
+auto kernel_dispatch_cb_records    = std::deque<kernel_dispatch_callback_record_t>{};
 
-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<rocprofiler_record_counter_t*>(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<std::mutex>{_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<rocprofiler_callback_tracing_kernel_dispatch_data_t*>(record.payload);
+
+        static auto _mutex = std::mutex{};
+        auto        _lk    = std::unique_lock<std::mutex>{_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<rocprofiler_buffer_tracing_hsa_api_record_t>{};
-auto marker_api_bf_records   = std::deque<rocprofiler_buffer_tracing_marker_api_record_t>{};
-auto hip_api_bf_records      = std::deque<rocprofiler_buffer_tracing_hip_api_record_t>{};
-auto kernel_dispatch_records = std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>{};
-auto memory_copy_records     = std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>{};
-auto scratch_memory_records  = std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>{};
+auto hsa_api_bf_records         = std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>{};
+auto marker_api_bf_records      = std::deque<rocprofiler_buffer_tracing_marker_api_record_t>{};
+auto hip_api_bf_records         = std::deque<rocprofiler_buffer_tracing_hip_api_record_t>{};
+auto kernel_dispatch_bf_records = std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>{};
+auto memory_copy_records        = std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>{};
+auto scratch_memory_records     = std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>{};
 auto corr_id_retire_records =
     std::deque<rocprofiler_buffer_tracing_correlation_id_retirement_record_t>{};
 
@@ -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<std::mutex>{_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<rocprofiler_buffer_tracing_kernel_dispatch_record_t*>(
                     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<rocprofiler_record_counter_t*>(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<std::string_view, rocprofiler_context_id_t*>{
     {"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_tracing_operation_t>(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<call_stack_t*>(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<std::string_view, std::string>{};
-        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<uint64_t>(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