AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

JSON schema + JSON node vs. Perfetto category consistency (#869)

Browse Source 
* Fix schema differences b/t JSON and perfetto

- "kernel_dispatches" (JSON) vs. "kernel_dispatch" (Perfetto) -> "kernel_dispatch"
- "scratch_api" (JSON) -> "scratch_memory"

* Update include/rocprofiler-sdk/cxx/serialization.hpp

- remove unnecessary includes causing circular dependency

* rocprofv3 docs

- Output formats
- JSON schema

* Spelling fix

* Update schema descriptions

* Improve assert log in test_perfetto_data

[ROCm/rocprofiler-sdk commit: b4dc4f8e92]
This commit is contained in:
Jonathan R. Madsen
2024-05-22 20:36:27 -05:00
committed by GitHub
parent 46a9637496
commit 9fde046078
18 changed files with 1746 additions and 70 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/source/docs/requirements.txt
+1
View File
@@ -1 +1,2 @@
doxysphinx
jsonschema2md
projects/rocprofiler-sdk/source/docs/rocprofv3-schema.json
+1397
View File
File diff suppressed because it is too large Load Diff
projects/rocprofiler-sdk/source/docs/rocprofv3.md
+285 -7
View File
@@ -1,4 +1,4 @@
# rocprofv3 user guide
# rocprofv3 User Guide
ROCProfiler SDK is a tooling infrastructure for profiling general-purpose GPU compute applications running on the ROCm platform. It supports application tracing to provide a big picture of the GPU application execution and kernel profiling to provide low-level hardware details from the performance counters.
@@ -106,7 +106,7 @@ ctest -R
## Usage
`rocprofv3` is a CLI tool that helps you quickly optimize applications and understand the low-level kernel details without requiring any modification in the source code. It is being developed to be backward compatible with its predecessor, `rocprof`, and to provide more features to help users profile their applications with better accuracy.
`rocprofv3` is a CLI tool that helps you quickly optimize applications and understand the low-level kernel details without requiring any modification in the source code. It is being developed to be backward compatible with its predecessor, `rocprof`, and to provide more features to help users profile their applications with better accuracy.
The following sections demonstrate the use of `rocprofv3` for application tracing and kernel profiling using various command-line options.
@@ -154,12 +154,12 @@ Below is the list of `rocprofv3` command-line options. Some options are used for
You can also see all the `rocprofv3` options using:
```bash
rocprofv3 --help
rocprofv3 --help
```
### Application tracing
Application tracing provides the big picture of a program’s execution by collecting data on the execution times of API calls and GPU commands, such as kernel execution, async memory copy, and barrier packets. This information can be used as the first step in the profiling process to answer important questions, such as how much percentage of time was spent on memory copy and which kernel took the longest time to execute.
Application tracing provides the big picture of a program’s execution by collecting data on the execution times of API calls and GPU commands, such as kernel execution, async memory copy, and barrier packets. This information can be used as the first step in the profiling process to answer important questions, such as how much percentage of time was spent on memory copy and which kernel took the longest time to execute.
To use `rocprofv3` for application tracing, run:
@@ -179,7 +179,7 @@ rocprofv3 --hip-trace < app_relative_path >
**Note: The tracing and counter colleciton options generates an additional agent info file. See** [Agent Info](#agent-info)
The above command generates a `hip_api_trace.csv` file prefixed with the process ID.
The above command generates a `hip_api_trace.csv` file prefixed with the process ID.
```bash
$ cat 238_hip_api_trace.csv
@@ -210,7 +210,7 @@ $ cat 208_hip_api_trace.csv
"HIP_COMPILER_API","__hipRegisterFatBinary",208,208,1,1508780270085955,1508780270096795
"HIP_COMPILER_API","__hipRegisterFunction",208,208,2,1508780270104242,1508780270115355
"HIP_COMPILER_API","__hipPushCallConfiguration",208,208,3,1508780613897816,1508780613898701
"HIP_COMPILER_API","__hipPopCallConfiguration",208,208,4,1508780613901714,1508780613902200
"HIP_COMPILER_API","__hipPopCallConfiguration",208,208,4,1508780613901714,1508780613902200
```
To describe the fields in the output file, see [Output file fields](#output-file-fields).
@@ -227,7 +227,7 @@ $ cat 238_agent_info.csv
#### HSA trace
The HIP runtime library is implemented with the low-level HSA runtime. HSA API tracing is more suited for advanced users who want to understand the application behavior at the lower level. In general, tracing at the HIP level is recommended for most users. You should use HSA trace only if you are familiar with HSA runtime.
The HIP runtime library is implemented with the low-level HSA runtime. HSA API tracing is more suited for advanced users who want to understand the application behavior at the lower level. In general, tracing at the HIP level is recommended for most users. You should use HSA trace only if you are familiar with HSA runtime.
HSA trace contains the start and end time of HSA runtime API calls and their asynchronous activities.
@@ -492,3 +492,281 @@ ctest -V
```bash
/usr/local/bin/python -m pip install -r requirements.txt
```
## Output Formats
rocprofv3 provides the following output formats:
- CSV (default)
- JSON
- PFTrace
Specification of the output format is via the `--output-format` command-line option. Format selection is case-insensitive
and multiple output formats can be selected via a single string concatenated by commas, colons, or semi-colons. Example:
`--output-format JSON` enables JSON output exclusively where as `--output-format csv,json,pftrace` enables all three output formats
for the run.
For trace visualization, use the PFTrace format and open the trace in [ui.perfetto.dev](https://ui.perfetto.dev/).
### JSON Output Schema
rocprofv3 supports a custom JSON output format designed for programmatic analysis. The schema is optimized for size
while factoring in usability. The JSON output can be generated via the `--output-format json` command-line option.
#### Properties
- **`rocprofiler-sdk-tool`** *(array)*: rocprofv3 data per process (each element represents a process).
- **Items** *(object)*: Data for rocprofv3.
- **`metadata`** *(object, required)*: Metadata related to the profiler session.
- **`pid`** *(integer, required)*: Process ID.
- **`init_time`** *(integer, required)*: Initialization time in nanoseconds.
- **`fini_time`** *(integer, required)*: Finalization time in nanoseconds.
- **`agents`** *(array, required)*: List of agents.
- **Items** *(object)*: Data for an agent.
- **`size`** *(integer, required)*: Size of the agent data.
- **`id`** *(object, required)*: Identifier for the agent.
- **`handle`** *(integer, required)*: Handle for the agent.
- **`type`** *(integer, required)*: Type of the agent.
- **`cpu_cores_count`** *(integer)*: Number of CPU cores.
- **`simd_count`** *(integer)*: Number of SIMD units.
- **`mem_banks_count`** *(integer)*: Number of memory banks.
- **`caches_count`** *(integer)*: Number of caches.
- **`io_links_count`** *(integer)*: Number of I/O links.
- **`cpu_core_id_base`** *(integer)*: Base ID for CPU cores.
- **`simd_id_base`** *(integer)*: Base ID for SIMD units.
- **`max_waves_per_simd`** *(integer)*: Maximum waves per SIMD.
- **`lds_size_in_kb`** *(integer)*: Size of LDS in KB.
- **`gds_size_in_kb`** *(integer)*: Size of GDS in KB.
- **`num_gws`** *(integer)*: Number of GWS (global work size).
- **`wave_front_size`** *(integer)*: Size of the wave front.
- **`num_xcc`** *(integer)*: Number of XCC (execution compute units).
- **`cu_count`** *(integer)*: Number of compute units (CUs).
- **`array_count`** *(integer)*: Number of arrays.
- **`num_shader_banks`** *(integer)*: Number of shader banks.
- **`simd_arrays_per_engine`** *(integer)*: SIMD arrays per engine.
- **`cu_per_simd_array`** *(integer)*: CUs per SIMD array.
- **`simd_per_cu`** *(integer)*: SIMDs per CU.
- **`max_slots_scratch_cu`** *(integer)*: Maximum slots for scratch CU.
- **`gfx_target_version`** *(integer)*: GFX target version.
- **`vendor_id`** *(integer)*: Vendor ID.
- **`device_id`** *(integer)*: Device ID.
- **`location_id`** *(integer)*: Location ID.
- **`domain`** *(integer)*: Domain identifier.
- **`drm_render_minor`** *(integer)*: DRM render minor version.
- **`num_sdma_engines`** *(integer)*: Number of SDMA engines.
- **`num_sdma_xgmi_engines`** *(integer)*: Number of SDMA XGMI engines.
- **`num_sdma_queues_per_engine`** *(integer)*: Number of SDMA queues per engine.
- **`num_cp_queues`** *(integer)*: Number of CP queues.
- **`max_engine_clk_ccompute`** *(integer)*: Maximum engine clock for compute.
- **`max_engine_clk_fcompute`** *(integer)*: Maximum engine clock for F compute.
- **`sdma_fw_version`** *(object)*: SDMA firmware version.
- **`uCodeSDMA`** *(integer, required)*: SDMA microcode version.
- **`uCodeRes`** *(integer, required)*: Reserved microcode version.
- **`fw_version`** *(object)*: Firmware version.
- **`uCode`** *(integer, required)*: Microcode version.
- **`Major`** *(integer, required)*: Major version.
- **`Minor`** *(integer, required)*: Minor version.
- **`Stepping`** *(integer, required)*: Stepping version.
- **`capability`** *(object, required)*: Agent capability flags.
- **`HotPluggable`** *(integer, required)*: Hot pluggable capability.
- **`HSAMMUPresent`** *(integer, required)*: HSAMMU present capability.
- **`SharedWithGraphics`** *(integer, required)*: Shared with graphics capability.
- **`QueueSizePowerOfTwo`** *(integer, required)*: Queue size is power of two.
- **`QueueSize32bit`** *(integer, required)*: Queue size is 32-bit.
- **`QueueIdleEvent`** *(integer, required)*: Queue idle event.
- **`VALimit`** *(integer, required)*: VA limit.
- **`WatchPointsSupported`** *(integer, required)*: Watch points supported.
- **`WatchPointsTotalBits`** *(integer, required)*: Total bits for watch points.
- **`DoorbellType`** *(integer, required)*: Doorbell type.
- **`AQLQueueDoubleMap`** *(integer, required)*: AQL queue double map.
- **`DebugTrapSupported`** *(integer, required)*: Debug trap supported.
- **`WaveLaunchTrapOverrideSupported`** *(integer, required)*: Wave launch trap override supported.
- **`WaveLaunchModeSupported`** *(integer, required)*: Wave launch mode supported.
- **`PreciseMemoryOperationsSupported`** *(integer, required)*: Precise memory operations supported.
- **`DEPRECATED_SRAM_EDCSupport`** *(integer, required)*: Deprecated SRAM EDC support.
- **`Mem_EDCSupport`** *(integer, required)*: Memory EDC support.
- **`RASEventNotify`** *(integer, required)*: RAS event notify.
- **`ASICRevision`** *(integer, required)*: ASIC revision.
- **`SRAM_EDCSupport`** *(integer, required)*: SRAM EDC support.
- **`SVMAPISupported`** *(integer, required)*: SVM API supported.
- **`CoherentHostAccess`** *(integer, required)*: Coherent host access.
- **`DebugSupportedFirmware`** *(integer, required)*: Debug supported firmware.
- **`Reserved`** *(integer, required)*: Reserved field.
- **`counters`** *(array, required)*: Array of counter objects.
- **Items** *(object)*
- **`agent_id`** *(object, required)*: Agent ID information.
- **`handle`** *(integer, required)*: Handle of the agent.
- **`id`** *(object, required)*: Counter ID information.
- **`handle`** *(integer, required)*: Handle of the counter.
- **`is_constant`** *(integer, required)*: Indicator if the counter value is constant.
- **`is_derived`** *(integer, required)*: Indicator if the counter value is derived.
- **`name`** *(string, required)*: Name of the counter.
- **`description`** *(string, required)*: Description of the counter.
- **`block`** *(string, required)*: Block information of the counter.
- **`expression`** *(string, required)*: Expression of the counter.
- **`dimension_ids`** *(array, required)*: Array of dimension IDs.
- **Items** *(integer)*: Dimension ID.
- **`strings`** *(object, required)*: String records.
- **`callback_records`** *(array)*: Callback records.
- **Items** *(object)*
- **`kind`** *(string, required)*: Kind of the record.
- **`operations`** *(array, required)*: Array of operations.
- **Items** *(string)*: Operation.
- **`buffer_records`** *(array)*: Buffer records.
- **Items** *(object)*
- **`kind`** *(string, required)*: Kind of the record.
- **`operations`** *(array, required)*: Array of operations.
- **Items** *(string)*: Operation.
- **`marker_api`** *(array)*: Marker API records.
- **Items** *(object)*
- **`key`** *(integer, required)*: Key of the record.
- **`value`** *(string, required)*: Value of the record.
- **`counters`** *(object)*: Counter records.
- **`dimension_ids`** *(array, required)*: Array of dimension IDs.
- **Items** *(object)*
- **`id`** *(integer, required)*: Dimension ID.
- **`instance_size`** *(integer, required)*: Size of the instance.
- **`name`** *(string, required)*: Name of the dimension.
- **`code_objects`** *(array, required)*: Code object records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the code object.
- **`code_object_id`** *(integer, required)*: ID of the code object.
- **`rocp_agent`** *(object, required)*: ROCP agent information.
- **`handle`** *(integer, required)*: Handle of the ROCP agent.
- **`hsa_agent`** *(object, required)*: HSA agent information.
- **`handle`** *(integer, required)*: Handle of the HSA agent.
- **`uri`** *(string, required)*: URI of the code object.
- **`load_base`** *(integer, required)*: Base address for loading.
- **`load_size`** *(integer, required)*: Size for loading.
- **`load_delta`** *(integer, required)*: Delta for loading.
- **`storage_type`** *(integer, required)*: Type of storage.
- **`memory_base`** *(integer, required)*: Base address for memory.
- **`memory_size`** *(integer, required)*: Size of memory.
- **`kernel_symbols`** *(array, required)*: Kernel symbol records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the kernel symbol.
- **`kernel_id`** *(integer, required)*: ID of the kernel.
- **`code_object_id`** *(integer, required)*: ID of the code object.
- **`kernel_name`** *(string, required)*: Name of the kernel.
- **`kernel_object`** *(integer, required)*: Object of the kernel.
- **`kernarg_segment_size`** *(integer, required)*: Size of the kernarg segment.
- **`kernarg_segment_alignment`** *(integer, required)*: Alignment of the kernarg segment.
- **`group_segment_size`** *(integer, required)*: Size of the group segment.
- **`private_segment_size`** *(integer, required)*: Size of the private segment.
- **`formatted_kernel_name`** *(string, required)*: Formatted name of the kernel.
- **`demangled_kernel_name`** *(string, required)*: Demangled name of the kernel.
- **`truncated_kernel_name`** *(string, required)*: Truncated name of the kernel.
- **`callback_records`** *(object, required)*: Callback record details.
- **`counter_collection`** *(array)*: Counter collection records.
- **Items** *(object)*
- **`dispatch_data`** *(object, required)*: Dispatch data details.
- **`size`** *(integer, required)*: Size of the dispatch data.
- **`correlation_id`** *(object, required)*: Correlation ID information.
- **`internal`** *(integer, required)*: Internal correlation ID.
- **`external`** *(integer, required)*: External correlation ID.
- **`dispatch_info`** *(object, required)*: Dispatch information details.
- **`size`** *(integer, required)*: Size of the dispatch information.
- **`agent_id`** *(object, required)*: Agent ID information.
- **`handle`** *(integer, required)*: Handle of the agent.
- **`queue_id`** *(object, required)*: Queue ID information.
- **`handle`** *(integer, required)*: Handle of the queue.
- **`kernel_id`** *(integer, required)*: ID of the kernel.
- **`dispatch_id`** *(integer, required)*: ID of the dispatch.
- **`private_segment_size`** *(integer, required)*: Size of the private segment.
- **`group_segment_size`** *(integer, required)*: Size of the group segment.
- **`workgroup_size`** *(object, required)*: Workgroup size information.
- **`x`** *(integer, required)*: X dimension.
- **`y`** *(integer, required)*: Y dimension.
- **`z`** *(integer, required)*: Z dimension.
- **`grid_size`** *(object, required)*: Grid size information.
- **`x`** *(integer, required)*: X dimension.
- **`y`** *(integer, required)*: Y dimension.
- **`z`** *(integer, required)*: Z dimension.
- **`records`** *(array, required)*: Records.
- **Items** *(object)*
- **`counter_id`** *(object, required)*: Counter ID information.
- **`handle`** *(integer, required)*: Handle of the counter.
- **`value`** *(number, required)*: Value of the counter.
- **`thread_id`** *(integer, required)*: Thread ID.
- **`arch_vgpr_count`** *(integer, required)*: Count of VGPRs.
- **`sgpr_count`** *(integer, required)*: Count of SGPRs.
- **`lds_block_size_v`** *(integer, required)*: Size of LDS block.
- **`buffer_records`** *(object, required)*: Buffer record details.
- **`kernel_dispatch`** *(array)*: Kernel dispatch records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the dispatch.
- **`kind`** *(integer, required)*: Kind of the dispatch.
- **`operation`** *(integer, required)*: Operation of the dispatch.
- **`thread_id`** *(integer, required)*: Thread ID.
- **`correlation_id`** *(object, required)*: Correlation ID information.
- **`internal`** *(integer, required)*: Internal correlation ID.
- **`external`** *(integer, required)*: External correlation ID.
- **`start_timestamp`** *(integer, required)*: Start timestamp.
- **`end_timestamp`** *(integer, required)*: End timestamp.
- **`dispatch_info`** *(object, required)*: Dispatch information details.
- **`size`** *(integer, required)*: Size of the dispatch information.
- **`agent_id`** *(object, required)*: Agent ID information.
- **`handle`** *(integer, required)*: Handle of the agent.
- **`queue_id`** *(object, required)*: Queue ID information.
- **`handle`** *(integer, required)*: Handle of the queue.
- **`kernel_id`** *(integer, required)*: ID of the kernel.
- **`dispatch_id`** *(integer, required)*: ID of the dispatch.
- **`private_segment_size`** *(integer, required)*: Size of the private segment.
- **`group_segment_size`** *(integer, required)*: Size of the group segment.
- **`workgroup_size`** *(object, required)*: Workgroup size information.
- **`x`** *(integer, required)*: X dimension.
- **`y`** *(integer, required)*: Y dimension.
- **`z`** *(integer, required)*: Z dimension.
- **`grid_size`** *(object, required)*: Grid size information.
- **`x`** *(integer, required)*: X dimension.
- **`y`** *(integer, required)*: Y dimension.
- **`z`** *(integer, required)*: Z dimension.
- **`hip_api`** *(array)*: HIP API records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the HIP API record.
- **`kind`** *(integer, required)*: Kind of the HIP API.
- **`operation`** *(integer, required)*: Operation of the HIP API.
- **`correlation_id`** *(object, required)*: Correlation ID information.
- **`internal`** *(integer, required)*: Internal correlation ID.
- **`external`** *(integer, required)*: External correlation ID.
- **`start_timestamp`** *(integer, required)*: Start timestamp.
- **`end_timestamp`** *(integer, required)*: End timestamp.
- **`thread_id`** *(integer, required)*: Thread ID.
- **`hsa_api`** *(array)*: HSA API records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the HSA API record.
- **`kind`** *(integer, required)*: Kind of the HSA API.
- **`operation`** *(integer, required)*: Operation of the HSA API.
- **`correlation_id`** *(object, required)*: Correlation ID information.
- **`internal`** *(integer, required)*: Internal correlation ID.
- **`external`** *(integer, required)*: External correlation ID.
- **`start_timestamp`** *(integer, required)*: Start timestamp.
- **`end_timestamp`** *(integer, required)*: End timestamp.
- **`thread_id`** *(integer, required)*: Thread ID.
- **`marker_api`** *(array)*: Marker (ROCTx) API records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the Marker API record.
- **`kind`** *(integer, required)*: Kind of the Marker API.
- **`operation`** *(integer, required)*: Operation of the Marker API.
- **`correlation_id`** *(object, required)*: Correlation ID information.
- **`internal`** *(integer, required)*: Internal correlation ID.
- **`external`** *(integer, required)*: External correlation ID.
- **`start_timestamp`** *(integer, required)*: Start timestamp.
- **`end_timestamp`** *(integer, required)*: End timestamp.
- **`thread_id`** *(integer, required)*: Thread ID.
- **`memory_copy`** *(array)*: Async memory copy records.
- **Items** *(object)*
- **`size`** *(integer, required)*: Size of the Marker API record.
- **`kind`** *(integer, required)*: Kind of the Marker API.
- **`operation`** *(integer, required)*: Operation of the Marker API.
- **`correlation_id`** *(object, required)*: Correlation ID information.
- **`internal`** *(integer, required)*: Internal correlation ID.
- **`external`** *(integer, required)*: External correlation ID.
- **`start_timestamp`** *(integer, required)*: Start timestamp.
- **`end_timestamp`** *(integer, required)*: End timestamp.
- **`thread_id`** *(integer, required)*: Thread ID.
- **`dst_agent_id`** *(object, required)*: Destination Agent ID.
- **`handle`** *(integer, required)*: Handle of the agent.
- **`src_agent_id`** *(object, required)*: Source Agent ID.
- **`handle`** *(integer, required)*: Handle of the agent.
- **`bytes`** *(integer, required)*: Bytes copied.
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/serialization.hpp
-2
View File
@@ -45,9 +45,7 @@
#include <cereal/types/functional.hpp>
#include <cereal/types/list.hpp>
#include <cereal/types/map.hpp>
#include <cereal/types/memory.hpp>
#include <cereal/types/optional.hpp>
#include <cereal/types/polymorphic.hpp>
#include <cereal/types/queue.hpp>
#include <cereal/types/set.hpp>
#include <cereal/types/stack.hpp>
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp
+4 -4
View File
@@ -45,7 +45,7 @@ write_json(tool_table* tool
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_deque,
std::deque<rocprofiler_tool_counter_collection_record_t>* counter_collection_deque,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_deque,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_api_deque)
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_memory_deque)
{
using JSONOutputArchive = cereal::MinimalJSONOutputArchive;
@@ -118,12 +118,12 @@ write_json(tool_table* tool
{
json_ar.setNextName("buffer_records");
json_ar.startNode();
json_ar(cereal::make_nvp("kernel_dispatches", *kernel_dispatch_deque));
json_ar(cereal::make_nvp("kernel_dispatch", *kernel_dispatch_deque));
json_ar(cereal::make_nvp("hip_api", *hip_api_deque));
json_ar(cereal::make_nvp("hsa_api", *hsa_api_deque));
json_ar(cereal::make_nvp("memory_copy", *memory_copy_deque));
json_ar(cereal::make_nvp("scratch_api", *scratch_api_deque));
json_ar(cereal::make_nvp("marker_api", *marker_api_deque));
json_ar(cereal::make_nvp("memory_copy", *memory_copy_deque));
json_ar(cereal::make_nvp("scratch_memory", *scratch_memory_deque));
json_ar.finishNode();
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp
+1 -1
View File
@@ -39,7 +39,7 @@ write_json(tool_table* tool
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_deque,
std::deque<rocprofiler_tool_counter_collection_record_t>* counter_collection_deque,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_deque,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_api_deque);
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_memory_deque);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp
+1 -1
View File
@@ -70,7 +70,7 @@ write_perfetto(
std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>* kernel_dispatch_data,
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* /*scratch_api_data*/)
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* /*scratch_memory_data*/)
{
namespace sdk = ::rocprofiler::sdk;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.hpp
+1 -1
View File
@@ -40,6 +40,6 @@ write_perfetto(
std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>* kernel_dispatch_data,
std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>* memory_copy_data,
std::deque<rocprofiler_buffer_tracing_marker_api_record_t>* marker_api_data,
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_api_data);
std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>* scratch_memory_data);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/tests/async-copy-tracing/validate.py
+12 -12
View File
@@ -37,11 +37,11 @@ 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("kernel_dispatch", sdk_data["callback_records"])
node_exists("memory_copies", sdk_data["callback_records"], 24)
node_exists("names", sdk_data["buffer_records"])
node_exists("kernel_dispatches", sdk_data["buffer_records"])
node_exists("kernel_dispatch", sdk_data["buffer_records"])
node_exists("memory_copies", sdk_data["buffer_records"], 12)
node_exists("hsa_api_traces", sdk_data["buffer_records"])
node_exists("hip_api_traces", sdk_data["buffer_records"], 0)
@@ -96,7 +96,7 @@ def test_timestamps(input_data):
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] <= itr["end_timestamp"]
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}"
assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}"
@@ -135,7 +135,7 @@ def test_internal_correlation_ids(input_data):
api_corr_ids_sorted = sorted(api_corr_ids)
api_corr_ids_unique = list(set(api_corr_ids))
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["correlation_id"]["internal"] in api_corr_ids_unique
for itr in sdk_data["buffer_records"]["memory_copies"]:
@@ -167,7 +167,7 @@ def test_external_correlation_ids(input_data):
assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}"
assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}"
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "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}"
@@ -196,10 +196,10 @@ def test_kernel_ids(input_data):
assert payload["kernel_id"] in symbol_info.keys()
assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()
@@ -207,17 +207,17 @@ 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"])
num_dispatches = len(sdk_data["buffer_records"]["kernel_dispatch"])
num_cb_dispatches = len(sdk_data["callback_records"]["kernel_dispatch"])
assert num_cb_dispatches == (3 * num_dispatches)
bf_seq_ids = []
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
bf_seq_ids.append(itr["dispatch_info"]["dispatch_id"])
cb_seq_ids = []
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
cb_seq_ids.append(itr["payload"]["dispatch_info"]["dispatch_id"])
bf_seq_ids = sorted(bf_seq_ids)
@@ -305,7 +305,7 @@ def test_retired_correlation_ids(input_data):
api_corr_ids[corr_id] = itr
async_corr_ids = {}
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
corr_id = itr["correlation_id"]["internal"]
assert corr_id not in async_corr_ids.keys()
projects/rocprofiler-sdk/tests/hip-graph-tracing/validate.py
+11 -11
View File
@@ -35,10 +35,10 @@ def test_data_structure(input_data):
node_exists("code_objects", sdk_data["callback_records"])
node_exists("kernel_symbols", sdk_data["callback_records"])
node_exists("hip_api_traces", sdk_data["callback_records"])
node_exists("kernel_dispatches", sdk_data["callback_records"])
node_exists("kernel_dispatch", sdk_data["callback_records"])
node_exists("names", sdk_data["buffer_records"])
node_exists("kernel_dispatches", sdk_data["buffer_records"])
node_exists("kernel_dispatch", sdk_data["buffer_records"])
node_exists("hip_api_traces", sdk_data["buffer_records"], 0)
@@ -89,7 +89,7 @@ def test_timestamps(input_data):
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] <= itr["end_timestamp"]
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}"
assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}"
@@ -128,7 +128,7 @@ def test_internal_correlation_ids(input_data):
api_corr_ids_sorted = sorted(api_corr_ids)
api_corr_ids_unique = list(set(api_corr_ids))
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["correlation_id"]["internal"] in api_corr_ids_unique
for itr in sdk_data["buffer_records"]["memory_copies"]:
@@ -160,7 +160,7 @@ def test_external_correlation_ids(input_data):
assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}"
assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}"
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "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}"
@@ -185,10 +185,10 @@ def test_kernel_ids(input_data):
assert payload["kernel_id"] in symbol_info.keys()
assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()
@@ -196,17 +196,17 @@ 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"])
num_dispatches = len(sdk_data["buffer_records"]["kernel_dispatch"])
num_cb_dispatches = len(sdk_data["callback_records"]["kernel_dispatch"])
assert num_cb_dispatches == (3 * num_dispatches)
bf_seq_ids = []
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
bf_seq_ids.append(itr["dispatch_info"]["dispatch_id"])
cb_seq_ids = []
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
cb_seq_ids.append(itr["payload"]["dispatch_info"]["dispatch_id"])
bf_seq_ids = sorted(bf_seq_ids)
projects/rocprofiler-sdk/tests/kernel-tracing/validate.py
+13 -13
View File
@@ -37,10 +37,10 @@ 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("kernel_dispatch", sdk_data["callback_records"])
node_exists("names", sdk_data["buffer_records"])
node_exists("kernel_dispatches", sdk_data["buffer_records"])
node_exists("kernel_dispatch", sdk_data["buffer_records"])
node_exists("memory_copies", sdk_data["buffer_records"], 0)
node_exists("hsa_api_traces", sdk_data["buffer_records"])
node_exists("hip_api_traces", sdk_data["buffer_records"], 0)
@@ -95,7 +95,7 @@ def test_timestamps(input_data):
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] <= itr["end_timestamp"]
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}"
assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}"
@@ -123,7 +123,7 @@ def test_total_runtime(input_data):
sdk_data = input_data["rocprofiler-sdk-json-tool"]
runtime_data = []
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
elapsed = itr["end_timestamp"] - itr["start_timestamp"]
runtime_data.append(elapsed) # in nanoseconds
@@ -148,7 +148,7 @@ def test_internal_correlation_ids(input_data):
api_corr_ids_sorted = sorted(api_corr_ids)
api_corr_ids_unique = list(set(api_corr_ids))
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["correlation_id"]["internal"] in api_corr_ids_unique
for itr in sdk_data["buffer_records"]["memory_copies"]:
@@ -174,7 +174,7 @@ def test_retired_correlation_ids(input_data):
api_corr_ids[corr_id] = itr
async_corr_ids = {}
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
corr_id = itr["correlation_id"]["internal"]
assert corr_id not in async_corr_ids.keys()
@@ -226,7 +226,7 @@ def test_external_correlation_ids(input_data):
assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}"
assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}"
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "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}"
@@ -251,10 +251,10 @@ def test_kernel_ids(input_data):
assert payload["kernel_id"] in symbol_info.keys()
assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()
@@ -262,17 +262,17 @@ 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"])
num_dispatches = len(sdk_data["buffer_records"]["kernel_dispatch"])
num_cb_dispatches = len(sdk_data["callback_records"]["kernel_dispatch"])
assert num_cb_dispatches == (3 * num_dispatches)
bf_seq_ids = []
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
bf_seq_ids.append(itr["dispatch_info"]["dispatch_id"])
cb_seq_ids = []
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
cb_seq_ids.append(itr["payload"]["dispatch_info"]["dispatch_id"])
bf_seq_ids = sorted(bf_seq_ids)
projects/rocprofiler-sdk/tests/page-migration/validate.py
+7 -7
View File
@@ -87,7 +87,7 @@ def test_data_structure(input_data):
node_exists("marker_api_traces", sdk_data["callback_records"], 0)
node_exists("names", sdk_data["buffer_records"])
node_exists("kernel_dispatches", sdk_data["buffer_records"])
node_exists("kernel_dispatch", sdk_data["buffer_records"])
node_exists("memory_copies", sdk_data["buffer_records"], 0)
node_exists("hsa_api_traces", sdk_data["buffer_records"])
node_exists("hip_api_traces", sdk_data["buffer_records"], 0)
@@ -117,7 +117,7 @@ def test_timestamps(input_data):
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] <= itr["end_timestamp"]
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] < itr["end_timestamp"]
assert itr["correlation_id"]["internal"] > 0
@@ -148,7 +148,7 @@ def test_internal_correlation_ids(input_data):
api_corr_ids_sorted = sorted(api_corr_ids)
api_corr_ids_unique = list(set(api_corr_ids))
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["correlation_id"]["internal"] in api_corr_ids_unique
for itr in sdk_data["buffer_records"]["memory_copies"]:
@@ -178,7 +178,7 @@ def test_external_correlation_ids(input_data):
assert itr["thread_id"] in extern_corr_ids
assert itr["correlation_id"]["external"] in extern_corr_ids
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["correlation_id"]["external"] > 0
assert itr["correlation_id"]["external"] in extern_corr_ids
@@ -206,10 +206,10 @@ def test_kernel_ids(input_data):
assert payload["kernel_id"] in symbol_info.keys()
assert payload["kernel_name"] == symbol_info[kern_id]["kernel_name"]
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["dispatch_info"]["kernel_id"] in symbol_info.keys()
for itr in sdk_data["callback_records"]["kernel_dispatches"]:
for itr in sdk_data["callback_records"]["kernel_dispatch"]:
assert itr["payload"]["dispatch_info"]["kernel_id"] in symbol_info.keys()
@@ -228,7 +228,7 @@ def test_retired_correlation_ids(input_data):
api_corr_ids[corr_id] = itr
async_corr_ids = {}
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
corr_id = itr["correlation_id"]["internal"]
assert corr_id not in async_corr_ids.keys()
projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py
+4 -2
View File
@@ -31,7 +31,7 @@ def test_perfetto_data(
"hip": ("hip_api", "hip_api"),
"hsa": ("hsa_api", "hsa_api"),
"marker": ("marker_api", "marker_api"),
"kernel": ("kernel_dispatch", "kernel_dispatches"),
"kernel": ("kernel_dispatch", "kernel_dispatch"),
"memory_copy": ("memory_copy", "memory_copy"),
}
@@ -45,4 +45,6 @@ def test_perfetto_data(
_pf_data = pftrace_data.loc[pftrace_data["category"] == pf_category]
_js_data = json_data["rocprofiler-sdk-tool"]["buffer_records"][js_category]
assert len(_pf_data) == len(_js_data)
assert len(_pf_data) == len(
_js_data
), f"{pf_category} ({len(_pf_data)}):\n\t{_pf_data}\n{js_category} ({len(_js_data)}):\n\t{_js_data}"
projects/rocprofiler-sdk/tests/rocprofv3/hsa-queue-dependency/validate.py
+1 -1
View File
@@ -70,7 +70,7 @@ def test_kernel_trace_json(json_data):
valid_kernel_names = ["copyA", "copyB", "copyC"]
buffer_records = data["buffer_records"]
buffer_names = data["strings"]["buffer_records"]
kernel_dispatch_data = buffer_records["kernel_dispatches"]
kernel_dispatch_data = buffer_records["kernel_dispatch"]
def get_kernel_name(kernel_id):
return data["kernel_symbols"][kernel_id]["formatted_kernel_name"]
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py
+1 -1
View File
@@ -304,7 +304,7 @@ def test_kernel_trace_json(json_data):
return data["kernel_symbols"][kernel_id]["formatted_kernel_name"]
kernels = []
for row in buffer_records["kernel_dispatches"]:
for row in buffer_records["kernel_dispatch"]:
dispatch_info = row["dispatch_info"]
kernel_name = get_kernel_name(dispatch_info["kernel_id"])
if re.search(r"__amd_rocclr_.*", kernel_name):
projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py
+1 -1
View File
@@ -134,7 +134,7 @@ def test_kernel_trace_json(json_data):
"_Z15matrixTransposePfS_i.kd",
"matrixTranspose(float*, float*, int)",
)
kernel_dispatch_data = data["buffer_records"]["kernel_dispatches"]
kernel_dispatch_data = data["buffer_records"]["kernel_dispatch"]
assert len(kernel_dispatch_data) == 1
for dispatch in kernel_dispatch_data:
dispatch_info = dispatch["dispatch_info"]
projects/rocprofiler-sdk/tests/scratch-memory-tracing/validate.py
+4 -4
View File
@@ -45,7 +45,7 @@ def test_data_structure(input_data):
node_exists("scratch_memory_traces", sdk_data["callback_records"], min_len=8)
node_exists("names", sdk_data["buffer_records"])
node_exists("kernel_dispatches", sdk_data["buffer_records"])
node_exists("kernel_dispatch", sdk_data["buffer_records"])
node_exists("memory_copies", sdk_data["buffer_records"], num_agents)
node_exists("hsa_api_traces", sdk_data["buffer_records"])
node_exists("hip_api_traces", sdk_data["buffer_records"], 0)
@@ -74,7 +74,7 @@ def test_timestamps(input_data):
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] <= itr["end_timestamp"]
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "memory_copies"]:
for itr in sdk_data["buffer_records"][titr]:
assert itr["start_timestamp"] < itr["end_timestamp"], f"[{titr}] {itr}"
assert itr["correlation_id"]["internal"] > 0, f"[{titr}] {itr}"
@@ -114,7 +114,7 @@ def test_internal_correlation_ids(input_data):
api_corr_ids_sorted = sorted(api_corr_ids)
api_corr_ids_unique = list(set(api_corr_ids))
for itr in sdk_data["buffer_records"]["kernel_dispatches"]:
for itr in sdk_data["buffer_records"]["kernel_dispatch"]:
assert itr["correlation_id"]["internal"] in api_corr_ids_unique
for itr in sdk_data["buffer_records"]["memory_copies"]:
@@ -146,7 +146,7 @@ def test_external_correlation_ids(input_data):
assert itr["thread_id"] in extern_corr_ids, f"[{titr}] {itr}"
assert itr["correlation_id"]["external"] in extern_corr_ids, f"[{titr}] {itr}"
for titr in ["kernel_dispatches", "memory_copies"]:
for titr in ["kernel_dispatch", "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}"
projects/rocprofiler-sdk/tests/tools/json-tool.cpp
+2 -2
View File
@@ -1475,7 +1475,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));
json_ar(cereal::make_nvp("kernel_dispatch", kernel_dispatch_cb_records));
json_ar(cereal::make_nvp("memory_copies", memory_copy_cb_records));
} catch(std::exception& e)
{
@@ -1490,7 +1490,7 @@ write_json(call_stack_t* _call_stack)
try
{
json_ar(cereal::make_nvp("names", buffer_names));
json_ar(cereal::make_nvp("kernel_dispatches", kernel_dispatch_bf_records));
json_ar(cereal::make_nvp("kernel_dispatch", kernel_dispatch_bf_records));
json_ar(cereal::make_nvp("memory_copies", memory_copy_bf_records));
json_ar(cereal::make_nvp("scratch_memory_traces", scratch_memory_records));
json_ar(cereal::make_nvp("page_migration", page_migration_records));