rocprofiler_kernel_dispatch_info_t + header record for buffered counter collection (#758)

* Update include/rocprofiler-sdk - defines.h - ROCPROFILER_VERSION_10_0 -> ROCPROFILER_SDK_VERSION_0_0 - fwd.h - rocprofiler_counter_record_kind_t - rocprofiler_kernel_dispatch_info_t - rocprofiler_record_counter_t - has dispatch id instead of correlation id - rocprofiler_counter_info_v0_t - added rocprofiler_counter_id_t field - added is_constant field - reordered better packing - dispatch_profile.h - added rocprofiler_profile_counting_dispatch_record_t for use as a header record for rocprofiler_profile_counting_dispatch_data_t - callback_tracing.h - rocprofiler_callback_tracing_kernel_dispatch_data_t uses rocprofiler_kernel_dispatch_info_t - buffer_tracing.h - rocprofiler_buffer_tracing_kernel_dispatch_record_t uses rocprofiler_kernel_dispatch_info_t * Update lib/rocprofiler-sdk/* - transition to rocprofiler_kernel_dispatch_info_t - set id and is_constant values for rocprofiler_counter_info_v0_t in rocprofiler_query_counter_info * Update lib/rocprofiler-sdk-tool - transition to rocprofiler_kernel_dispatch_info_t * Update lib/rocprofiler-sdk/counters/tests/core.cpp - transition to rocprofiler_kernel_dispatch_info_t * Update samples - transition to rocprofiler_kernel_dispatch_info_t - transition to rocprofiler_counter_record_kind_t * Update tests - transition to rocprofiler_kernel_dispatch_info_t - transition to rocprofiler_counter_record_kind_t - improve integration test validation for counter-collection - update serialization for new/additional types * Fix tests/counter-collection/validate.py - loosen restrictions on the length of counter description * Update include/rocprofiler-sdk/buffer_tracing.h - remove accidental packed attribute * Update lib/rocprofiler-sdk/counters/xml/derived_counters.xml - Add description for TCC_TAG_STALL_sum (reference: https://rocm.docs.amd.com/en/develop/conceptual/gpu-arch/mi300-mi200-performance-counters.html) * Update tests/page-migration/validate.py [ROCm/rocprofiler-sdk commit: 07537b6231]
2024-04-12 17:30:34 -05:00
parent 4f99edbad5
commit 2aef3c3d15
23 changed files with 489 additions and 273 deletions
@@ -194,10 +194,10 @@ def test_kernel_ids(input_data):
            assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]

    for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
-        assert itr["kernel_id"] in symbol_info.keys()
+        assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()

    for itr in sdk_data["callback_records"]["kernel_dispatches"]:
-        assert itr["payload"]["kernel_id"] in symbol_info.keys()
+        assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()


 def test_kernel_dispatch_ids(input_data):
@@ -211,11 +211,11 @@ def test_kernel_dispatch_ids(input_data):

    bf_seq_ids = []
    for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
-        bf_seq_ids.append(itr["dispatch_id"])
+        bf_seq_ids.append(itr["dispatch_info"]["dispatch_id"])

    cb_seq_ids = []
    for itr in sdk_data["callback_records"]["kernel_dispatches"]:
-        cb_seq_ids.append(itr["payload"]["dispatch_id"])
+        cb_seq_ids.append(itr["payload"]["dispatch_info"]["dispatch_id"])

    bf_seq_ids = sorted(bf_seq_ids)
    cb_seq_ids = sorted(cb_seq_ids)
@@ -97,6 +97,13 @@ save(ArchiveT& ar, rocprofiler_queue_id_t data)
    SAVE_DATA_FIELD(handle);
 }

+template <typename ArchiveT>
+void
+save(ArchiveT& ar, rocprofiler_counter_id_t data)
+{
+    SAVE_DATA_FIELD(handle);
+}
+
 template <typename ArchiveT>
 void
 save(ArchiveT& ar, rocprofiler_correlation_id_t data)
@@ -270,11 +277,9 @@ save(ArchiveT& ar, rocprofiler_callback_tracing_scratch_memory_data_t data)

 template <typename ArchiveT>
 void
-save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data)
+save(ArchiveT& ar, rocprofiler_kernel_dispatch_info_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);
@@ -285,20 +290,33 @@ save(ArchiveT& ar, rocprofiler_callback_tracing_kernel_dispatch_data_t data)
    SAVE_DATA_FIELD(group_segment_size);
 }

+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(dispatch_info);
+}
+
 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);
+    SAVE_DATA_FIELD(dispatch_info);
+}
+
+template <typename ArchiveT>
+void
+save(ArchiveT& ar, rocprofiler_profile_counting_dispatch_record_t data)
+{
+    SAVE_DATA_FIELD(size);
+    SAVE_DATA_FIELD(num_records);
+    SAVE_DATA_FIELD(correlation_id);
+    SAVE_DATA_FIELD(dispatch_info);
 }

 template <typename ArchiveT>
@@ -339,7 +357,7 @@ save(ArchiveT& ar, rocprofiler_record_counter_t data)
 {
    SAVE_DATA_FIELD(id);
    SAVE_DATA_FIELD(counter_value);
-    SAVE_DATA_FIELD(correlation_id);
+    SAVE_DATA_FIELD(dispatch_id);
 }

 template <typename ArchiveT>
@@ -367,14 +385,7 @@ save(ArchiveT& ar, rocprofiler_buffer_tracing_kernel_dispatch_record_t data)
    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);
-    SAVE_DATA_FIELD(grid_size);
+    SAVE_DATA_FIELD(dispatch_info);
 }

 template <typename ArchiveT>
@@ -705,6 +716,19 @@ save(ArchiveT& ar, const rocprofiler_agent_t& data)
    generate("caches", data.caches, data.caches_count);
    generate("io_links", data.io_links, data.io_links_count);
 }
+
+template <typename ArchiveT>
+void
+save(ArchiveT& ar, rocprofiler_counter_info_v0_t data)
+{
+    SAVE_DATA_FIELD(id);
+    SAVE_DATA_BITFIELD("is_constant", is_constant);
+    SAVE_DATA_BITFIELD("is_derived", is_derived);
+    SAVE_DATA_CSTR(name);
+    SAVE_DATA_CSTR(description);
+    SAVE_DATA_CSTR(block);
+    SAVE_DATA_CSTR(expression);
+}
 }  // namespace cereal

 #undef SAVE_DATA_FIELD
@@ -22,24 +22,50 @@ def test_data_structure(input_data):


 def test_counter_values(input_data):
-    data = input_data
-    agent_data = data["rocprofiler-sdk-json-tool"]["agents"]
-    counter_data = data["rocprofiler-sdk-json-tool"]["buffer_records"][
-        "counter_collection"
-    ]
+    data = input_data["rocprofiler-sdk-json-tool"]
+    agent_data = data["agents"]
+    counter_info = data["counter_info"]
+    counter_data = data["buffer_records"]["counter_collection"]

-    scaling_factor = 1
-    for itr in agent_data:
-        if itr["type"] == 2 and itr["wave_front_size"] > 0:
-            scaling_factor = 64 / itr["wave_front_size"]
-            break
+    for itr in counter_info:
+        if itr["is_constant"] == 1 and itr["name"] == "size":
+            continue
+        assert itr["id"]["handle"] > 0, f"{itr}"
+        assert itr["is_constant"] in (0, 1), f"{itr}"
+        assert itr["is_derived"] in (0, 1), f"{itr}"
+        assert len(itr["name"]) >= 4, f"{itr}"
+        assert len(itr["description"]) >= 4, f"{itr}"
+        if itr["is_constant"] == 0:
+            if itr["is_derived"] == 0:
+                assert len(itr["block"]) > 0, f"{itr}"
+            if itr["is_derived"] == 1:
+                assert len(itr["expression"]) > 0, f"{itr}"
+
+    def get_agent(agent_id):
+        for itr in agent_data:
+            if itr["id"]["handle"] == agent_id["handle"]:
+                return itr
+        return None
+
+    def get_scaling_factor(agent_id):
+        agent = get_agent(agent_id)
+        assert agent is not None, f"id={agent_id}"
+        if agent["type"] == 2 and agent["wave_front_size"] > 0:
+            return 64 / agent["wave_front_size"]
+        return 0

    for itr in counter_data:
-        value = itr["counter_value"]
-        if int(round(value, 0)) > 0:
-            assert int(round(value, 0)) == int(
-                round(1 * scaling_factor, 0)
-            ), f"agent_data:\n{agent_data}\n\ncounter_data:\n{counter_data}"
+        assert itr["num_records"] == len(itr["records"]), f"itr={itr}"
+        agent_id = itr["dispatch_info"]["agent_id"]
+        agent = get_agent(agent_id)
+        scaling_factor = get_scaling_factor(agent_id)
+        assert agent is not None, f"itr={itr}\nagent={agent}"
+        for ritr in itr["records"]:
+            value = ritr["counter_value"]
+            if int(round(value, 0)) > 0:
+                assert int(round(value, 0)) == int(
+                    round(1 * scaling_factor, 0)
+                ), f"itr={itr}\nagent={agent}"


 if __name__ == "__main__":
@@ -254,10 +254,10 @@ def test_kernel_ids(input_data):
            assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]

    for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
-        assert itr["kernel_id"] in symbol_info.keys()
+        assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()

    for itr in sdk_data["callback_records"]["kernel_dispatches"]:
-        assert itr["payload"]["kernel_id"] in symbol_info.keys()
+        assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()


 def test_kernel_dispatch_ids(input_data):
@@ -271,11 +271,11 @@ def test_kernel_dispatch_ids(input_data):

    bf_seq_ids = []
    for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
-        bf_seq_ids.append(itr["dispatch_id"])
+        bf_seq_ids.append(itr["dispatch_info"]["dispatch_id"])

    cb_seq_ids = []
    for itr in sdk_data["callback_records"]["kernel_dispatches"]:
-        cb_seq_ids.append(itr["payload"]["dispatch_id"])
+        cb_seq_ids.append(itr["payload"]["dispatch_info"]["dispatch_id"])

    bf_seq_ids = sorted(bf_seq_ids)
    cb_seq_ids = sorted(cb_seq_ids)
@@ -209,7 +209,10 @@ def test_kernel_ids(input_data):
            assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]

    for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
-        assert itr["kernel_id"] in symbol_info.keys()
+        assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()
+
+    for itr in sdk_data["callback_records"]["kernel_dispatches"]:
+        assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()


 def test_retired_correlation_ids(input_data):
@@ -40,6 +40,8 @@
 #include <rocprofiler-sdk/buffer.h>
 #include <rocprofiler-sdk/buffer_tracing.h>
 #include <rocprofiler-sdk/callback_tracing.h>
+#include <rocprofiler-sdk/counters.h>
+#include <rocprofiler-sdk/dispatch_profile.h>
 #include <rocprofiler-sdk/external_correlation.h>
 #include <rocprofiler-sdk/fwd.h>
 #include <rocprofiler-sdk/internal_threading.h>
@@ -62,6 +64,7 @@
 #include <map>
 #include <mutex>
 #include <shared_mutex>
+#include <stdexcept>
 #include <string>
 #include <string_view>
 #include <thread>
@@ -549,11 +552,59 @@ struct scratch_memory_callback_record_t
    }
 };

+struct profile_counting_record
+{
+    profile_counting_record(rocprofiler_profile_counting_dispatch_record_t hdr)
+    : header{hdr}
+    {}
+
+    rocprofiler_profile_counting_dispatch_record_t header = {};
+    std::vector<rocprofiler_record_counter_t>      data   = {};
+
+    profile_counting_record()                                   = default;
+    ~profile_counting_record()                                  = default;
+    profile_counting_record(const profile_counting_record&)     = default;
+    profile_counting_record(profile_counting_record&&) noexcept = default;
+    profile_counting_record& operator=(const profile_counting_record&) = default;
+    profile_counting_record& operator=(profile_counting_record&&) noexcept = default;
+
+    template <typename ArchiveT>
+    void save(ArchiveT& ar) const
+    {
+        cereal::save(ar, header);
+        auto _data = data;
+        for(auto& itr : _data)
+        {
+            auto _counter_id = rocprofiler_counter_id_t{};
+            ROCPROFILER_CALL(rocprofiler_query_record_counter_id(itr.id, &_counter_id),
+                             "failed to query counter id");
+            itr.id = _counter_id.handle;
+        }
+
+        ar(cereal::make_nvp("records", _data));
+    }
+
+    void emplace_back(rocprofiler_record_counter_t val)
+    {
+        if(*this != val)
+            throw std::runtime_error{"invalid profile_counting_record::emplace_back(...)"};
+        data.emplace_back(val);
+    }
+
+    bool operator==(rocprofiler_record_counter_t rhs) const
+    {
+        return (header.dispatch_info.dispatch_id == rhs.dispatch_id);
+    }
+
+    bool operator!=(rocprofiler_record_counter_t rhs) const { return !(*this == rhs); }
+};
+
+auto counter_info                  = std::deque<rocprofiler_counter_info_v0_t>{};
 auto code_object_records           = std::deque<code_object_callback_record_t>{};
 auto kernel_symbol_records         = std::deque<kernel_symbol_callback_record_t>{};
 auto hsa_api_cb_records            = std::deque<hsa_api_callback_record_t>{};
 auto marker_api_cb_records         = std::deque<marker_api_callback_record_t>{};
-auto counter_collection_bf_records = std::deque<rocprofiler_record_counter_t>{};
+auto counter_collection_bf_records = std::deque<profile_counting_record>{};
 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>{};
@@ -583,7 +634,8 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data,
    static std::unordered_map<uint64_t, rocprofiler_profile_config_id_t> profile_cache = {};

    auto search_cache = [&]() {
-        if(auto pos = profile_cache.find(dispatch_data.agent_id.handle); pos != profile_cache.end())
+        if(auto pos = profile_cache.find(dispatch_data.dispatch_info.agent_id.handle);
+           pos != profile_cache.end())
        {
            *config = pos->second;
            return true;
@@ -610,7 +662,7 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data,

    // Iterate through the agents and get the counters available on that agent
    ROCPROFILER_CALL(rocprofiler_iterate_agent_supported_counters(
-                         dispatch_data.agent_id,
+                         dispatch_data.dispatch_info.agent_id,
                         []([[maybe_unused]] rocprofiler_agent_id_t id,
                            rocprofiler_counter_id_t*               counters,
                            size_t                                  num_counters,
@@ -626,18 +678,30 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data,
                         static_cast<void*>(&gpu_counters)),
                     "Could not fetch supported counters");

+    for(auto& counter : gpu_counters)
+    {
+        auto info = rocprofiler_counter_info_v0_t{};
+
+        ROCPROFILER_CALL(
+            rocprofiler_query_counter_info(
+                counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast<void*>(&info)),
+            "Could not query counter_id");
+
+        counter_info.emplace_back(info);
+    }
+
    std::vector<rocprofiler_counter_id_t> collect_counters;
    // Look for the counters contained in counters_to_collect in gpu_counters
    for(auto& counter : gpu_counters)
    {
-        rocprofiler_counter_info_v0_t version;
+        rocprofiler_counter_info_v0_t info;

        ROCPROFILER_CALL(
            rocprofiler_query_counter_info(
-                counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast<void*>(&version)),
+                counter, ROCPROFILER_COUNTER_INFO_VERSION_0, static_cast<void*>(&info)),
            "Could not query counter_id");

-        if(counters_to_collect.count(std::string(version.name)) > 0)
+        if(counters_to_collect.count(std::string(info.name)) > 0)
        {
            collect_counters.push_back(counter);
        }
@@ -645,12 +709,13 @@ dispatch_callback(rocprofiler_profile_counting_dispatch_data_t dispatch_data,

    // Create a colleciton profile for the counters
    rocprofiler_profile_config_id_t profile;
-    ROCPROFILER_CALL(
-        rocprofiler_create_profile_config(
-            dispatch_data.agent_id, collect_counters.data(), collect_counters.size(), &profile),
-        "Could not construct profile cfg");
+    ROCPROFILER_CALL(rocprofiler_create_profile_config(dispatch_data.dispatch_info.agent_id,
+                                                       collect_counters.data(),
+                                                       collect_counters.size(),
+                                                       &profile),
+                     "Could not construct profile cfg");

-    profile_cache.emplace(dispatch_data.agent_id.handle, profile);
+    profile_cache.emplace(dispatch_data.dispatch_info.agent_id.handle, profile);
    // Return the profile to collect those counters for this dispatch
    *config = profile;
 }
@@ -865,10 +930,21 @@ 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)
+        else if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS &&
+                header->kind == ROCPROFILER_COUNTER_RECORD_PROFILE_COUNTING_DISPATCH_HEADER)
+        {
+            auto* profiler_record =
+                static_cast<rocprofiler_profile_counting_dispatch_record_t*>(header->payload);
+            counter_collection_bf_records.emplace_back(*profiler_record);
+        }
+        else if(header->category == ROCPROFILER_BUFFER_CATEGORY_COUNTERS &&
+                header->kind == ROCPROFILER_COUNTER_RECORD_VALUE)
        {
            auto* profiler_record = static_cast<rocprofiler_record_counter_t*>(header->payload);
-            counter_collection_bf_records.emplace_back(*profiler_record);
+            if(counter_collection_bf_records.empty())
+                throw std::runtime_error{
+                    "missing rocprofiler_profile_counting_dispatch_record_t (header)"};
+            counter_collection_bf_records.back().emplace_back(*profiler_record);
        }
        else
        {
@@ -1427,7 +1503,8 @@ tool_fini(void* tool_data)
              << ", 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_value_records=" << counter_collection_bf_records.size()
+              << "...\n"
              << std::flush;

    auto* _call_stack = static_cast<call_stack_t*>(tool_data);
@@ -1501,6 +1578,7 @@ write_json(call_stack_t* _call_stack)
        json_ar.finishNode();

        json_ar(cereal::make_nvp("agents", agents));
+        json_ar(cereal::make_nvp("counter_info", counter_info));
        if(_call_stack) json_ar(cereal::make_nvp("call_stack", *_call_stack));

        json_ar.setNextName("callback_records");
@@ -1611,7 +1689,8 @@ write_perfetto()
        }

        for(auto itr : kernel_dispatch_bf_records)
-            agent_queue_ids[itr.agent_id.handle].emplace(itr.queue_id.handle);
+            agent_queue_ids[itr.dispatch_info.agent_id.handle].emplace(
+                itr.dispatch_info.queue_id.handle);
    }

    auto thread_tracks = std::unordered_map<rocprofiler_thread_id_t, ::perfetto::Track>{};
@@ -1804,10 +1883,11 @@ write_perfetto()
        auto demangled = std::unordered_map<std::string_view, std::string>{};
        for(auto itr : kernel_dispatch_bf_records)
        {
-            const kernel_symbol_callback_record_t* sym = nullptr;
+            const auto&                            info = itr.dispatch_info;
+            const kernel_symbol_callback_record_t* sym  = nullptr;
            for(const auto& kitr : kernel_symbol_records)
            {
-                if(kitr.payload.kernel_id == itr.kernel_id)
+                if(kitr.payload.kernel_id == info.kernel_id)
                {
                    sym = &kitr;
                    break;
@@ -1815,7 +1895,7 @@ write_perfetto()
            }

            auto  name  = std::string_view{sym->payload.kernel_name};
-            auto& track = agent_queue_tracks.at(itr.agent_id.handle).at(itr.queue_id.handle);
+            auto& track = agent_queue_tracks.at(info.agent_id.handle).at(info.queue_id.handle);

            if(demangled.find(name) == demangled.end())
            {
@@ -1833,21 +1913,21 @@ write_perfetto()
                "kind",
                itr.kind,
                "agent",
-                itr.agent_id.handle,
+                info.agent_id.handle,
                "corr_id",
                itr.correlation_id.internal,
                "queue",
-                itr.queue_id.handle,
+                info.queue_id.handle,
                "kernel_id",
-                itr.kernel_id,
+                info.kernel_id,
                "private_segment_size",
-                itr.private_segment_size,
+                info.private_segment_size,
                "group_segment_size",
-                itr.group_segment_size,
+                info.group_segment_size,
                "workgroup_size",
-                itr.workgroup_size.x * itr.workgroup_size.y * itr.workgroup_size.z,
+                info.workgroup_size.x * info.workgroup_size.y * info.workgroup_size.z,
                "grid_size",
-                itr.grid_size.x * itr.grid_size.y * itr.grid_size.z);
+                info.grid_size.x * info.grid_size.y * info.grid_size.z);

            TRACE_EVENT_END(rocprofiler::trait::name<rocprofiler::category::kernel_dispatch>::value,
                            track,