From 6743608258bd81c3c15973666f48eb7ffd06c846 Mon Sep 17 00:00:00 2001 From: "Jonathan R. Madsen" Date: Mon, 9 Sep 2024 11:20:55 -0500 Subject: [PATCH] rocprofv3: summary reports + more JSON metadata (#1029) * Move include/rocprofiler-sdk/cxx/details/delimit.hpp to tokenize.hpp * Update docs/how-to/using-rocprofv3.rst - fix code block indents - reorder rocprofv3 options, limit them to important options - add docs for `--runtime-trace` * Update rocprofv3.py - parser argument groups - new `--runtime-trace` option - new `--summary` option - new `--summary-per-domain` option - new `--summary-groups` option - new `--summary-output-file` option - new `--summary-units` option * Update lib/rocprofiler-sdk/hsa/async_copy.cpp - fix async copy operation names: add "MEMORY_COPY_" prefix * lib/rocprofiler-sdk-tool: update statistics.{hpp,cpp} - statistics<>::get_percent function - stats_entry_t struct - stats_formatter struct - percentage struct - std::to_string(::rocprofiler::tool::percentage) * lib/rocprofiler-sdk-tool: update domain_type.{hpp,cpp} - reorder domain_type enum values * lib/rocprofiler-sdk-tool: update generateCSV.{hpp,cpp} - separate writing CSV from accumulating statistics - a lot of functionality was moved to statistics.{hpp,cpp} * lib/rocprofiler-sdk-tool: update output_file.{hpp,cpp} - output_stream_t struct - get_output_stream(...) returns output_stream_t instance * lib/rocprofiler-sdk-tool: update generateJSON.cpp - update get_output_stream usage to output_stream_t * lib/rocprofiler-sdk-tool: update generateOTF2.cpp - header include order tweak * lib/rocprofiler-sdk-tool: update buffered_output.hpp - stats_data_t was renamed to stats_entry_t * lib/rocprofiler-sdk-tool: update generatePerfetto.cpp - header include tweak * lib/rocprofiler-sdk-tool: update tmp_file_buffer.hpp - emit warning message if write_ring_buffer fails after offloading instead of aborting - prefer placement new instead of assignment in write_ring_buffer * lib/rocprofiler-sdk-tool: add generateStats.{hpp,cpp} - functions for accumulating statistics * Update tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt - accommodate tweak to CSV output file name for HIP and HSA traces * lib/rocprofiler-sdk-tool: update config.{hpp,cpp} - new config variables - stats_summary - stats_summary_per_domain - summary_output - stats_summary_unit_value - stats_summary_unit - stats_summary_file - stats_summary_groups - support output keys for hostname: %hostname% / %h * lib/rocprofiler-sdk-tool: update tool.cpp - support summary output * Documentation fixes * Test for summary output * Update tests/bin/transpose to use more ROCTx - also support building with the roctracer ROCTx * Remove roctxMark from OTF2 + fix kernel-rename tests - following more ROCTx calls in transpose, kernel-rename validation had to be updated * JSON metadata + JSON summary - add serialization support for config - add serialization support for statistics - additions to json spec - rocprofiler-sdk-tool/metadata/config - rocprofiler-sdk-tool/metadata/command - rocprofiler-sdk-tool/summary - config output_keys support for NVIDIA %q{} syntax - config output_keys support keys within keys * rocprofv3 --summary-groups warning if no domain matches - emit warning if a regex in for summary groups did not match any domain names * Compile fix for lib/rocprofiler-sdk-tool/tool.cpp - get_config().scratch_memory_trace - pass contributions to write_json * Update rocprofv3.py to preload rocprofiler-sdk-roctx - appended to LD_PRELOAD when args.marker_trace is enabled * Fix ReST link errors about subtitle underline being too short * Patch tokenization of config::stats_summary_groups - guard against array values of empty strings * Tweak rocprofv3 summary test - input-summary.yaml (used by rocprofv3-test-summary-inp-yaml-execute) only provides one summary group regex * Disable LD_PRELOAD of librocprofiler-sdk-roctx.so - this causes problems in the sanitizers, will be addressed in another PR [ROCm/rocprofiler-sdk commit: 395f01b6895784c357d8c5753cbeadc0ddf3be35] --- .../rocprofiler-sdk/source/bin/rocprofv3.py | 367 ++++++++++++------ .../data/{hip_stats.csv => hip_api_stats.csv} | 0 .../data/{hsa_trace.csv => hsa_api_trace.csv} | 0 .../source/docs/how-to/using-rocprofv3.rst | 354 +++++++++-------- .../cxx/details/CMakeLists.txt | 2 +- .../cxx/details/{delimit.hpp => tokenize.hpp} | 58 +++ .../lib/rocprofiler-sdk-tool/CMakeLists.txt | 3 + .../rocprofiler-sdk-tool/buffered_output.hpp | 2 +- .../lib/rocprofiler-sdk-tool/config.cpp | 107 ++++- .../lib/rocprofiler-sdk-tool/config.hpp | 120 +++++- .../lib/rocprofiler-sdk-tool/domain_type.cpp | 10 +- .../lib/rocprofiler-sdk-tool/domain_type.hpp | 6 +- .../lib/rocprofiler-sdk-tool/generateCSV.cpp | 231 ++++------- .../lib/rocprofiler-sdk-tool/generateCSV.hpp | 65 ++-- .../lib/rocprofiler-sdk-tool/generateJSON.cpp | 41 +- .../lib/rocprofiler-sdk-tool/generateJSON.hpp | 2 + .../lib/rocprofiler-sdk-tool/generateOTF2.cpp | 9 +- .../rocprofiler-sdk-tool/generatePerfetto.cpp | 17 +- .../rocprofiler-sdk-tool/generateStats.cpp | 350 +++++++++++++++++ .../rocprofiler-sdk-tool/generateStats.hpp | 63 +++ .../lib/rocprofiler-sdk-tool/output_file.cpp | 19 +- .../lib/rocprofiler-sdk-tool/output_file.hpp | 57 ++- .../lib/rocprofiler-sdk-tool/statistics.cpp | 75 ++++ .../lib/rocprofiler-sdk-tool/statistics.hpp | 120 ++++++ .../rocprofiler-sdk-tool/tmp_file_buffer.hpp | 53 ++- .../source/lib/rocprofiler-sdk-tool/tool.cpp | 33 +- .../lib/rocprofiler-sdk/hsa/async_copy.cpp | 2 +- .../tests/bin/transpose/CMakeLists.txt | 67 +++- .../tests/bin/transpose/transpose.cpp | 42 +- .../tests/pytest-packages/tests/rocprofv3.py | 24 ++ .../tests/rocprofv3/CMakeLists.txt | 1 + .../tests/rocprofv3/kernel-rename/validate.py | 20 +- .../tests/rocprofv3/summary/CMakeLists.txt | 127 ++++++ .../tests/rocprofv3/summary/conftest.py | 108 ++++++ .../rocprofv3/summary/input-summary.yaml | 11 + .../tests/rocprofv3/summary/pytest.ini | 5 + .../tests/rocprofv3/summary/validate.py | 267 +++++++++++++ .../tracing-hip-in-libraries/CMakeLists.txt | 10 +- .../tracing-hip-in-libraries/validate.py | 9 +- .../tests/rocprofv3/tracing/validate.py | 16 +- 40 files changed, 2293 insertions(+), 580 deletions(-) rename projects/rocprofiler-sdk/source/docs/data/{hip_stats.csv => hip_api_stats.csv} (100%) rename projects/rocprofiler-sdk/source/docs/data/{hsa_trace.csv => hsa_api_trace.csv} (100%) rename projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/{delimit.hpp => tokenize.hpp} (75%) create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.cpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.hpp create mode 100644 projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.cpp create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/summary/CMakeLists.txt create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/summary/conftest.py create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/summary/input-summary.yaml create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/summary/pytest.ini create mode 100644 projects/rocprofiler-sdk/tests/rocprofv3/summary/validate.py diff --git a/projects/rocprofiler-sdk/source/bin/rocprofv3.py b/projects/rocprofiler-sdk/source/bin/rocprofv3.py index fadaa32b3d..cf84fa96a8 100755 --- a/projects/rocprofiler-sdk/source/bin/rocprofv3.py +++ b/projects/rocprofiler-sdk/source/bin/rocprofv3.py @@ -83,8 +83,8 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins formatter_class=argparse.RawTextHelpFormatter, ) - def add_parser_bool_argument(*args, **kwargs): - parser.add_argument( + def add_parser_bool_argument(gparser, *args, **kwargs): + gparser.add_argument( *args, **kwargs, action=booleanArgAction, @@ -95,90 +95,15 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins metavar="BOOL", ) - # Add the arguments - add_parser_bool_argument( - "--hip-trace", - help="For Collecting HIP Traces (runtime + compiler)", - ) - add_parser_bool_argument( - "--hip-runtime-trace", - help="For Collecting HIP Runtime API Traces", - ) - add_parser_bool_argument( - "--hip-compiler-trace", - help="For Collecting HIP Compiler generated code Traces", - ) - add_parser_bool_argument( - "--marker-trace", - help="For Collecting Marker (ROCTx) Traces", - ) - add_parser_bool_argument( - "--kernel-trace", - help="For Collecting Kernel Dispatch Traces", - ) - add_parser_bool_argument( - "--memory-copy-trace", - help="For Collecting Memory Copy Traces", - ) - add_parser_bool_argument( - "--scratch-memory-trace", - help="For Collecting Scratch Memory operations Traces", - ) - add_parser_bool_argument( - "--stats", - help="For Collecting statistics of enabled tracing types", - ) - add_parser_bool_argument( - "--hsa-trace", - help="For Collecting HSA Traces (core + amd + image + finalizer)", - ) - add_parser_bool_argument( - "--hsa-core-trace", - help="For Collecting HSA API Traces (core API)", - ) - add_parser_bool_argument( - "--hsa-amd-trace", - help="For Collecting HSA API Traces (AMD-extension API)", - ) - add_parser_bool_argument( - "--hsa-image-trace", - help="For Collecting HSA API Traces (Image-extenson API)", - ) - add_parser_bool_argument( - "--hsa-finalizer-trace", - help="For Collecting HSA API Traces (Finalizer-extension API)", - ) - add_parser_bool_argument( - "-s", - "--sys-trace", - help="For Collecting HIP, HSA, Marker (ROCTx), Memory copy, Scratch memory, and Kernel dispatch traces", - ) - add_parser_bool_argument( - "-M", - "--mangled-kernels", - help="Do not demangle the kernel names", - ) - add_parser_bool_argument( - "-T", - "--truncate-kernels", - help="Truncate the demangled kernel names", - ) - add_parser_bool_argument( - "-L", - "--list-metrics", - help="List metrics for counter collection", - ) - add_parser_bool_argument( - "--kernel-rename", - help="Use region names defined by roctxRangePush/roctxRangePop regions to rename the kernels", - ) - parser.add_argument( + io_options = parser.add_argument_group("I/O options") + + io_options.add_argument( "-i", "--input", help="Input file for counter collection", required=False, ) - parser.add_argument( + io_options.add_argument( "-o", "--output-file", help="For the output file name", @@ -186,7 +111,7 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins type=str, required=False, ) - parser.add_argument( + io_options.add_argument( "-d", "--output-directory", help="For adding output path where the output files will be saved", @@ -194,7 +119,7 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins type=str, required=False, ) - parser.add_argument( + io_options.add_argument( "--output-format", help="For adding output format (supported formats: csv, json, pftrace, otf2)", nargs="+", @@ -202,41 +127,188 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins choices=("csv", "json", "pftrace", "otf2"), type=str.lower, ) - parser.add_argument( + io_options.add_argument( "--log-level", help="Set the log level", default=None, choices=("fatal", "error", "warning", "info", "trace", "env"), type=str.lower, ) - parser.add_argument( + + aggregate_tracing_options = parser.add_argument_group("Aggregate tracing options") + + add_parser_bool_argument( + aggregate_tracing_options, + "-r", + "--runtime-trace", + help="Collect tracing data for HIP runtime API, Marker (ROCTx) API, Memory operations (copies and scratch), and Kernel dispatches. Similar to --sys-trace but without tracing HIP compiler API and the underlying HSA API.", + ) + add_parser_bool_argument( + aggregate_tracing_options, + "-s", + "--sys-trace", + help="Collect tracing data for HIP API, HSA API, Marker (ROCTx) API, Memory operations (copies and scratch), and Kernel dispatches.", + ) + + basic_tracing_options = parser.add_argument_group("Basic tracing options") + + # Add the arguments + add_parser_bool_argument( + basic_tracing_options, + "--hip-trace", + help="For collecting HIP Traces (runtime + compiler)", + ) + add_parser_bool_argument( + basic_tracing_options, + "--marker-trace", + help="For collecting Marker (ROCTx) Traces", + ) + add_parser_bool_argument( + basic_tracing_options, + "--kernel-trace", + help="For collecting Kernel Dispatch Traces", + ) + add_parser_bool_argument( + basic_tracing_options, + "--memory-copy-trace", + help="For collecting Memory Copy Traces", + ) + add_parser_bool_argument( + basic_tracing_options, + "--scratch-memory-trace", + help="For collecting Scratch Memory operations Traces", + ) + add_parser_bool_argument( + basic_tracing_options, + "--hsa-trace", + help="For collecting HSA Traces (core + amd + image + finalizer)", + ) + + extended_tracing_options = parser.add_argument_group("Granular tracing options") + + add_parser_bool_argument( + extended_tracing_options, + "--hip-runtime-trace", + help="For collecting HIP Runtime API Traces, e.g. public HIP API functions starting with 'hip' (i.e. hipSetDevice).", + ) + add_parser_bool_argument( + extended_tracing_options, + "--hip-compiler-trace", + help="For collecting HIP Compiler generated code Traces, e.g. HIP API functions starting with '__hip' (i.e. __hipRegisterFatBinary).", + ) + add_parser_bool_argument( + extended_tracing_options, + "--hsa-core-trace", + help="For collecting HSA API Traces (core API), e.g. HSA functions prefixed with only 'hsa_' (i.e. hsa_init).", + ) + add_parser_bool_argument( + extended_tracing_options, + "--hsa-amd-trace", + help="For collecting HSA API Traces (AMD-extension API), e.g. HSA function prefixed with 'hsa_amd_' (i.e. hsa_amd_coherency_get_type).", + ) + add_parser_bool_argument( + extended_tracing_options, + "--hsa-image-trace", + help="For collecting HSA API Traces (Image-extenson API), e.g. HSA functions prefixed with only 'hsa_ext_image_' (i.e. hsa_ext_image_get_capability).", + ) + add_parser_bool_argument( + extended_tracing_options, + "--hsa-finalizer-trace", + help="For collecting HSA API Traces (Finalizer-extension API), e.g. HSA functions prefixed with only 'hsa_ext_program_' (i.e. hsa_ext_program_create).", + ) + + post_processing_options = parser.add_argument_group("Post-processing tracing options") + + add_parser_bool_argument( + post_processing_options, + "--stats", + help="For collecting statistics of enabled tracing types", + ) + add_parser_bool_argument( + post_processing_options, + "-S", + "--summary", + help="Output single summary of tracing data at the conclusion of the profiling session", + ) + add_parser_bool_argument( + post_processing_options, + "-D", + "--summary-per-domain", + help="Output summary for each tracing domain at the conclusion of the profiling session", + ) + post_processing_options.add_argument( + "--summary-groups", + help="Output a summary for each set of domains matching the regular expression, e.g. 'KERNEL_DISPATCH|MEMORY_COPY' will generate a summary from all the tracing data in the KERNEL_DISPATCH and MEMORY_COPY domains; '*._API' will generate a summary from all the tracing data in the HIP_API, HSA_API, and MARKER_API domains", + nargs="+", + default=None, + type=str, + metavar="REGULAR_EXPRESSION", + ) + + summary_options = parser.add_argument_group("Summary options") + + summary_options.add_argument( + "--summary-output-file", + help="Output summary to a file, stdout, or stderr (default: stderr)", + default=None, + type=str, + ) + summary_options.add_argument( + "-u", + "--summary-units", + help="Timing units for output summary", + default=None, + type=str, + choices=("sec", "msec", "usec", "nsec"), + ) + + kernel_naming_options = parser.add_argument_group("Kernel naming options") + + add_parser_bool_argument( + kernel_naming_options, + "-M", + "--mangled-kernels", + help="Do not demangle the kernel names", + ) + add_parser_bool_argument( + kernel_naming_options, + "-T", + "--truncate-kernels", + help="Truncate the demangled kernel names", + ) + add_parser_bool_argument( + kernel_naming_options, + "--kernel-rename", + help="Use region names defined by roctxRangePush/roctxRangePop regions to rename the kernels", + ) + + filter_options = parser.add_argument_group("Filtering options") + + filter_options.add_argument( "--kernel-include-regex", - help="Include the kernels matching this filter", + help="Include the kernels matching this filter from counter-collection and thread-trace data (non-matching kernels will be excluded)", default=None, type=str, metavar="REGULAR_EXPRESSION", ) - parser.add_argument( + filter_options.add_argument( "--kernel-exclude-regex", - help="Exclude the kernels matching this filter", + help="Exclude the kernels matching this filter from counter-collection and thread-trace data (applied after --kernel-include-regex option)", default=None, type=str, metavar="REGULAR_EXPRESSION", ) - parser.add_argument( + filter_options.add_argument( "--kernel-iteration-range", help="Iteration range", nargs="+", default=None, type=str, ) - parser.add_argument( - "--preload", - help="Libraries to prepend to LD_PRELOAD (usually for sanitizers)", - default=os.environ.get("ROCPROF_PRELOAD", "").split(":"), - nargs="*", - ) - parser.add_argument( + + perfetto_options = parser.add_argument_group("Perfetto-specific options") + + perfetto_options.add_argument( "--perfetto-backend", help="Perfetto data collection backend. 'system' mode requires starting traced and perfetto daemons", default=None, @@ -244,21 +316,21 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins nargs=1, choices=("inprocess", "system"), ) - parser.add_argument( + perfetto_options.add_argument( "--perfetto-buffer-size", help="Size of buffer for perfetto output in KB. default: 1 GB", default=None, type=int, metavar="KB", ) - parser.add_argument( + perfetto_options.add_argument( "--perfetto-buffer-fill-policy", help="Policy for handling new records when perfetto has reached the buffer limit", default=None, type=str, choices=("discard", "ring_buffer"), ) - parser.add_argument( + perfetto_options.add_argument( "--perfetto-shmem-size-hint", help="Perfetto shared memory size hint in KB. default: 64 KB", default=None, @@ -266,6 +338,24 @@ For MPI applications (or other job launchers such as SLURM), place rocprofv3 ins metavar="KB", ) + display_options = parser.add_argument_group("Display options") + + add_parser_bool_argument( + display_options, + "-L", + "--list-metrics", + help="List metrics for counter collection", + ) + + advanced_options = parser.add_argument_group("Advanced options") + + advanced_options.add_argument( + "--preload", + help="Libraries to prepend to LD_PRELOAD (useful for sanitizer libraries)", + default=os.environ.get("ROCPROF_PRELOAD", "").split(":"), + nargs="*", + ) + if args is None: args = sys.argv[1:] @@ -488,9 +578,13 @@ def run(app_args, args, **kwargs): if not _overwrite: pass elif _prepend: - app_env[env_var] = "{}{}{}".format(_val, _join_char, _curr_val) + app_env[env_var] = ( + "{}{}{}".format(_val, _join_char, _curr_val) if _val else _curr_val + ) elif _append: - app_env[env_var] = "{}{}{}".format(_curr_val, _join_char, _val) + app_env[env_var] = ( + "{}{}{}".format(_curr_val, _join_char, _val) if _val else _curr_val + ) elif _overwrite: _write_env_value() else: @@ -504,12 +598,13 @@ def run(app_args, args, **kwargs): ROCM_DIR = os.path.dirname(ROCPROFV3_DIR) ROCPROF_TOOL_LIBRARY = f"{ROCM_DIR}/lib/rocprofiler-sdk/librocprofiler-sdk-tool.so" ROCPROF_SDK_LIBRARY = f"{ROCM_DIR}/lib/librocprofiler-sdk.so" + ROCPROF_ROCTX_LIBRARY = f"{ROCM_DIR}/lib/librocprofiler-sdk-roctx.so" - args.preload = [itr for itr in args.preload if itr] - if args.preload: - update_env("LD_PRELOAD", ":".join(args.preload), prepend=True) + prepend_preload = [itr for itr in args.preload if itr] + append_preload = [ROCPROF_TOOL_LIBRARY, ROCPROF_SDK_LIBRARY] - update_env("LD_PRELOAD", f"{ROCPROF_TOOL_LIBRARY}:{ROCPROF_SDK_LIBRARY}", append=True) + update_env("LD_PRELOAD", ":".join(prepend_preload), prepend=True) + update_env("LD_PRELOAD", ":".join(append_preload), append=True) update_env( "ROCP_TOOL_LIBRARIES", @@ -555,6 +650,16 @@ def run(app_args, args, **kwargs): ): setattr(args, itr, True) + if args.runtime_trace: + for itr in ( + "hip_runtime_trace", + "marker_trace", + "kernel_trace", + "memory_copy_trace", + "scratch_memory_trace", + ): + setattr(args, itr, True) + if args.hip_trace: for itr in ("compiler", "runtime"): setattr(args, f"hip_{itr}_trace", True) @@ -584,14 +689,54 @@ def run(app_args, args, **kwargs): trace_count += 1 if val else 0 trace_opts += ["--{}".format(opt.replace("_", "-"))] - if trace_count == 0 and args.stats: - fatal_error( - "No tracing options were enabled for --stats option. Tracing options:\n\t{}".format( - "\n\t".join(trace_opts) - ) - ) + # if marker tracing was requested, LD_PRELOAD the rocprofiler-sdk-roctx library + # to override the roctx symbols of an app linked to the old roctracer roctx + # if args.marker_trace: + # update_env("LD_PRELOAD", ROCPROF_ROCTX_LIBRARY, append=True) + + if trace_count == 0: + # if no tracing was enabled but the options below were enabled, raise an error + for oitr in [ + "stats", + "summary", + "summary-per-domain", + "summary-groups", + "summary-output-file", + "summary-units", + ]: + _attr = oitr.replace("-", "_") + if not hasattr(args, _attr): + fatal_error( + f"Internal error. parser does not support --{oitr} argument (i.e. args.{_attr})" + ) + elif getattr(args, _attr): + _len = max([len(f"{key}") for key in args.keys()]) + _args = "\n\t".join( + sorted([f"{key:<{_len}} = {val}" for key, val in args.items()]) + ) + fatal_error( + "No tracing options were enabled for --{} option.\nConfiguration argument values:\n\t{}\nTracing options:\n\t{}".format( + oitr, f"{_args}", "\n\t".join(trace_opts) + ) + ) + + _summary_groups = "##@@##".join(args.summary_groups) if args.summary_groups else None + _summary_output_fname = args.summary_output_file + if _summary_output_fname is None: + _summary_output_fname = "stderr" + elif _summary_output_fname.lower() in ("stdout", "stderr"): + _summary_output_fname = _summary_output_fname.lower() update_env("ROCPROF_STATS", args.stats, overwrite_if_true=True) + update_env("ROCPROF_STATS_SUMMARY", args.summary, overwrite_if_true=True) + update_env("ROCPROF_STATS_SUMMARY_UNITS", args.summary_units, overwrite=True) + update_env("ROCPROF_STATS_SUMMARY_OUTPUT", _summary_output_fname, overwrite=True) + update_env("ROCPROF_STATS_SUMMARY_GROUPS", _summary_groups, overwrite=True) + update_env( + "ROCPROF_STATS_SUMMARY_PER_DOMAIN", + args.summary_per_domain, + overwrite_if_true=True, + ) update_env( "ROCPROF_DEMANGLE_KERNELS", not args.mangled_kernels, diff --git a/projects/rocprofiler-sdk/source/docs/data/hip_stats.csv b/projects/rocprofiler-sdk/source/docs/data/hip_api_stats.csv similarity index 100% rename from projects/rocprofiler-sdk/source/docs/data/hip_stats.csv rename to projects/rocprofiler-sdk/source/docs/data/hip_api_stats.csv diff --git a/projects/rocprofiler-sdk/source/docs/data/hsa_trace.csv b/projects/rocprofiler-sdk/source/docs/data/hsa_api_trace.csv similarity index 100% rename from projects/rocprofiler-sdk/source/docs/data/hsa_trace.csv rename to projects/rocprofiler-sdk/source/docs/data/hsa_api_trace.csv diff --git a/projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst b/projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst index 2d7a2e1772..2a6f7a30a3 100644 --- a/projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst +++ b/projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst @@ -17,19 +17,19 @@ The following sections demonstrate the use of ``rocprofv3`` for application trac .. code-block:: bash - export PATH=$PATH:/opt/rocm/bin + export PATH=$PATH:/opt/rocm/bin Before you start tracing or profiling your HIP application using ``rocprofv3``, build the application using: .. code-block:: bash - cmake -B -DCMAKE_PREFIX_PATH=/opt/rocm - cmake --build --target all --parallel + cmake -B -DCMAKE_PREFIX_PATH=/opt/rocm + cmake --build --target all --parallel Options --------- -Here is the list of ``rocprofv3`` command-line options. Some options are used for application tracing and some for kernel profiling while the output control options control the presentation and redirection of the generated output. +Here is the sample of commonly used ``rocprofv3`` command-line options. Some options are used for application tracing and some for kernel profiling while the output control options control the presentation and redirection of the generated output. .. list-table:: rocprofv3 options :header-rows: 1 @@ -38,44 +38,32 @@ Here is the list of ``rocprofv3`` command-line options. Some options are used fo - Description - Use + * - ``-i`` \| ``--input`` + - Specifies the input file. JSON and YAML formats support configuration of all command-line options whereas the text format only supports specifying HW counters. + - Run Configuration + + * - ``-d`` \| ``--output-directory`` + - Specifies the path for the output files. Supports special keys: ``%hostname%``, ``%pid%``, ``%rank%``, etc. + - Output control + + * - ``-o`` \| ``--output-file`` + - Specifies the name of the output file. Note that this name is appended to the default names (_api_trace or counter_collection.csv) of the generated files'. Supports special keys: ``%hostname%``, ``%pid%``, ``%rank%``, etc. + - Output control + + * - ``--output-format`` + - For adding output format (supported formats: csv, json, pftrace) + - Output control + + * - ``-r`` \| ``--runtime-trace`` + - Collects HIP (runtime), memory copy, marker, scratch memory, and kernel dispatch traces. + - Application Tracing + + * - ``-s`` \| ``--sys-trace`` + - Collects HIP, HSA, memory copy, marker, scratch memory, and kernel dispatch traces. + - Application Tracing + * - ``--hip-trace`` - - Collects HIP runtime traces. - - Application tracing - - * - ``--hip-runtime-trace`` - - Collects HIP runtime API traces. - - Application tracing - - * - ``--hip-compiler-trace`` - - Collects HIP compiler-generated code traces. - - Application tracing - - * - ``--scratch-memory-trace`` - - Collects scratch memory operations traces. - - Application tracing - - * - ``--hsa-trace`` - - Collects HSA API traces. - - Application tracing - - * - ``--hsa-core-trace`` - - Collects HSA API traces (core API). - - Application tracing - - * - ``--hsa-amd-trace`` - - Collects HSA API traces (AMD-extension API). - - Application tracing - - * - ``--hsa-image-trace`` - - Collects HSA API Traces (Image-extension API). - - Application tracing - - * - ``--hsa-finalizer-trace`` - - Collects HSA API traces (Finalizer-extension API). - - Application tracing - - * - ``--stats`` - - For Collecting statistics of enabled tracing types + - Collects HIP runtime and compiler traces. - Application tracing * - ``--kernel-trace`` @@ -90,37 +78,49 @@ Here is the list of ``rocprofv3`` command-line options. Some options are used fo - Collects memory copy traces. - Application tracing - * - ``--sys-trace`` - - Collects HIP, HSA, memory copy, marker, and kernel dispatch traces. - - Application Tracing + * - ``--scratch-memory-trace`` + - Collects scratch memory operations traces. + - Application tracing - * - ``-i`` - - Specifies the input file. - - Kernel profiling with text file. Tracing and profiling with JSON and YAML. + * - ``--hsa-trace`` + - Collects HSA API traces. + - Application tracing + + * - ``--hip-runtime-trace`` + - Collects HIP runtime API traces. + - Application tracing + + * - ``--hsa-core-trace`` + - Collects HSA API traces (core API). + - Application tracing + + * - ``--hsa-amd-trace`` + - Collects HSA API traces (AMD-extension API). + - Application tracing + + * - ``--stats`` + - For Collecting statistics of enabled tracing types + - Application tracing + + * - ``-p`` \| ``--summary`` + - Display summary of collected data + - Application tracing * - ``--kernel-include-regex`` - Include the kernels matching this filter. - - Kernel profiling + - Kernel Dispatch Counter Collection * - ``--kernel-exclude-regex`` - Exclude the kernels matching this filter. - - Kernel profiling + - Kernel Dispatch Counter Collection * - ``--kernel-iteration-range`` - Iteration range for each kernel that match the filter [start-stop]. - - Kernel profiling + - Kernel Dispatch Counter Collection * - ``-L`` \| ``--list-metrics`` - List metrics for counter collection. - - Kernel profiling - - * - ``-d`` \| ``--output-directory`` - - Specifies the path for the output files. - - Output control - - * - ``-o`` \| ``--output-file`` - - Specifies the name of the output file. Note that this name is appended to the default names (_api_trace or counter_collection.csv) of the generated files'. - - Output control + - Kernel Dispatch Counter Collection * - ``-M`` \| ``--mangled-kernels`` - Overrides the default demangling of kernel names. @@ -130,19 +130,11 @@ Here is the list of ``rocprofv3`` command-line options. Some options are used fo - Truncates the demangled kernel names for improved readability. - Output control - * - ``--output-format`` - - For adding output format (supported formats: csv, json, pftrace) - - Output control - - * - ``--preload`` - - Libraries to prepend to LD_PRELOAD (usually for sanitizers) - - Extension - -You can also see all the ``rocprofv3`` options using: +You can see the full list of all the ``rocprofv3`` options using: .. code-block:: bash - rocprofv3 --help + rocprofv3 --help Application tracing --------------------- @@ -153,7 +145,7 @@ To use ``rocprofv3`` for application tracing, run: .. code-block:: bash - rocprofv3 -- + rocprofv3 -- HIP trace +++++++++++ @@ -164,15 +156,15 @@ To trace HIP runtime APIs, use: .. code-block:: bash - rocprofv3 --hip-trace -- < app_relative_path > + rocprofv3 --hip-trace -- < app_relative_path > -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. .. code-block:: shell - $ cat 238_hip_api_trace.csv + $ cat 238_hip_api_trace.csv -Here are the contents of `hip_api_trace.csv` file: +Here are the contents of ``hip_api_trace.csv`` file: .. csv-table:: HIP runtime api trace :file: /data/hip_compile_trace.csv @@ -183,15 +175,15 @@ To trace HIP compile time APIs, use: .. code-block:: shell - rocprofv3 --hip-compiler-trace -- < app_relative_path > + rocprofv3 --hip-compiler-trace -- < app_relative_path > -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. .. code-block:: shell - $ cat 208_hip_api_trace.csv + $ cat 208_hip_api_trace.csv -Here are the contents of `hip_api_trace.csv` file: +Here are the contents of ``hip_api_trace.csv`` file: .. csv-table:: HIP compile time api trace :file: /data/hip_compile_trace.csv @@ -209,18 +201,18 @@ HSA trace contains the start and end time of HSA runtime API calls and their asy .. code-block:: bash - rocprofv3 --hsa-trace -- < app_relative_path > + rocprofv3 --hsa-trace -- < app_relative_path > -The above command generates a `hsa_api_trace.csv` file prefixed with process ID. Note that the contents of this file have been truncated for demonstration purposes. +The above command generates a ``hsa_api_trace.csv`` file prefixed with process ID. Note that the contents of this file have been truncated for demonstration purposes. .. code-block:: shell - $ cat 197_hsa_api_trace.csv + $ cat 197_hsa_api_trace.csv -Here are the contents of `hsa_api_trace.csv` file: +Here are the contents of ``hsa_api_trace.csv`` file: .. csv-table:: HSA api trace - :file: /data/hsa_trace.csv + :file: /data/hsa_api_trace.csv :widths: 10,10,10,10,10,20,20 :header-rows: 1 @@ -231,7 +223,7 @@ Marker trace In certain situations, such as debugging performance issues in large-scale GPU programs, API-level tracing might be too fine-grained to provide a big picture of the program execution. In such cases, it is helpful to define specific tasks to be traced. -To specify the tasks for tracing, enclose the respective source code with the API calls provided by the `ROCTX` library. This process is also known as instrumentation. As the scope of code for instrumentation is defined using the enclosing API calls, it is called a range. A range is a programmer-defined task that has a well-defined start and end code scope. You can also refine the scope specified within a range using further nested ranges. ``rocprofv3`` also reports the timelines for these nested ranges. +To specify the tasks for tracing, enclose the respective source code with the API calls provided by the ``ROCTX`` library. This process is also known as instrumentation. As the scope of code for instrumentation is defined using the enclosing API calls, it is called a range. A range is a programmer-defined task that has a well-defined start and end code scope. You can also refine the scope specified within a range using further nested ranges. ``rocprofv3`` also reports the timelines for these nested ranges. Here is a list of useful APIs for code instrumentation. @@ -241,41 +233,41 @@ Here is a list of useful APIs for code instrumentation. - ``roctxRangePop``: Stops the current nested range. - ``roctxRangeStop``: Stops the given range. -See how to use `rocTX` APIs in the MatrixTranspose application below: +See how to use ``rocTX`` APIs in the MatrixTranspose application below: .. code-block:: bash - roctxMark("before hipLaunchKernel"); - int rangeId = roctxRangeStart("hipLaunchKernel range"); - roctxRangePush("hipLaunchKernel"); + roctxMark("before hipLaunchKernel"); + int rangeId = roctxRangeStart("hipLaunchKernel range"); + roctxRangePush("hipLaunchKernel"); - // Launching kernel from host - hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0,0,gpuTransposeMatrix,gpuMatrix, WIDTH); + // Launching kernel from host + hipLaunchKernelGGL(matrixTranspose, dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y), dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y), 0,0,gpuTransposeMatrix,gpuMatrix, WIDTH); - roctxMark("after hipLaunchKernel"); + roctxMark("after hipLaunchKernel"); - // Memory transfer from device to host - roctxRangePush("hipMemcpy"); + // Memory transfer from device to host + roctxRangePush("hipMemcpy"); - hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); + hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM * sizeof(float), hipMemcpyDeviceToHost); - roctxRangePop(); // for "hipMemcpy" - roctxRangePop(); // for "hipLaunchKernel" - roctxRangeStop(rangeId); + roctxRangePop(); // for "hipMemcpy" + roctxRangePop(); // for "hipLaunchKernel" + roctxRangeStop(rangeId); To trace the API calls enclosed within the range, use: .. code-block:: bash - rocprofv3 --marker-trace -- < app_relative_path > + rocprofv3 --marker-trace -- < app_relative_path > -Running the preceding command generates a `marker_api_trace.csv` file prefixed with the process ID. +Running the preceding command generates a ``marker_api_trace.csv`` file prefixed with the process ID. .. code-block:: shell - $ cat 210_marker_api_trace.csv + $ cat 210_marker_api_trace.csv -Here are the contents of `marker_api_trace.csv` file: +Here are the contents of ``marker_api_trace.csv`` file: .. csv-table:: Marker api trace :file: /data/marker_api_trace.csv @@ -291,15 +283,15 @@ To trace kernel dispatch traces, use: .. code-block:: shell - rocprofv3 --kernel-trace -- < app_relative_path > + rocprofv3 --kernel-trace -- < app_relative_path > -The above command generates a `kernel_trace.csv` file prefixed with the process ID. +The above command generates a ``kernel_trace.csv`` file prefixed with the process ID. .. code-block:: shell - $ cat 199_kernel_trace.csv + $ cat 199_kernel_trace.csv -Here are the contents of `kernel_trace.csv` file: +Here are the contents of ``kernel_trace.csv`` file: .. csv-table:: Kernel trace :file: /data/kernel_trace.csv @@ -315,15 +307,15 @@ To trace memory moves across the application, use: .. code-block:: shell - rocprofv3 –-memory-copy-trace -- < app_relative_path > + rocprofv3 –-memory-copy-trace -- < app_relative_path > -The above command generates a `memory_copy_trace.csv` file prefixed with the process ID. +The above command generates a ``memory_copy_trace.csv`` file prefixed with the process ID. .. code-block:: shell - $ cat 197_memory_copy_trace.csv + $ cat 197_memory_copy_trace.csv -Here are the contents of `memory_copy_trace.csv` file: +Here are the contents of ``memory_copy_trace.csv`` file: .. csv-table:: Memory copy trace :file: /data/memory_copy_trace.csv @@ -332,16 +324,36 @@ Here are the contents of `memory_copy_trace.csv` file: For the description of the fields in the output file, see :ref:`output-file-fields`. -Sys trace -+++++++++++ +Runtime trace ++++++++++++++++ + +This is a short-hand option which attempts to target the most relevant tracing options for a standard user by +excluding tracing the HSA runtime API and HIP compiler API. + +The HSA runtime API is excluded because it is a lower-level API upon which HIP and OpenMP target are built and +thus, tends to be an implementation detail not relevant to most users. The HIP compiler API is excluded +because these are functions which are automatically inserted during HIP compilation and thus, also tend to be +implementation details which are not relevant to most users. + +At present, `--runtime-trace` enables tracing the HIP runtime API, the marker API, kernel dispatches, and +memory operations (copies and scratch). + +.. code-block:: shell + + rocprofv3 –-runtime-trace -- < app_relative_path > + +Running the above command generates ``hip_api_trace.csv``, ``kernel_trace.csv``, ``memory_copy_trace.csv``, ``scratch_memory_trace.csv``,and ``marker_api_trace.csv`` (if ``rocTX`` APIs are specified in the application) files prefixed with the process ID. + +System trace +++++++++++++++ This is an all-inclusive option to collect all the above-mentioned traces. .. code-block:: shell - rocprofv3 –-sys-trace -- < app_relative_path > + rocprofv3 –-sys-trace -- < app_relative_path > -Running the above command generates `hip_api_trace.csv`, `hsa_api_trace.csv`, `kernel_trace.csv`, `memory_copy_trace.csv`, and `marker_api_trace.csv` (if `rocTX` APIs are specified in the application) files prefixed with the process ID. +Running the above command generates ``hip_api_trace.csv``, ``hsa_api_trace.csv``, ``kernel_trace.csv``, ``memory_copy_trace.csv``, and ``marker_api_trace.csv`` (if ``rocTX`` APIs are specified in the application) files prefixed with the process ID. Scratch memory trace ++++++++++++++++++++++ @@ -350,7 +362,7 @@ This option collects scratch memory operation's traces. Scratch is an address sp .. code-block:: shell - rocprofv3 --scratch-memory-trace -- < app_relative_path > + rocprofv3 --scratch-memory-trace -- < app_relative_path > Stats ++++++++ @@ -360,18 +372,18 @@ A higher percentage in statistics can help user focus on the API/function that h .. code-block:: shell - rocprofv3 --stats --hip-trace -- < app_relative_path > + rocprofv3 --stats --hip-trace -- < app_relative_path > -The above command generates a `hip_stats.csv` and `hip_api_trace` file prefixed with the process ID. +The above command generates a ``hip_api_stats.csv`` and ``hip_api_trace`` file prefixed with the process ID. .. code-block:: shell - $ cat hip_stats.csv + $ cat hip_api_stats.csv -Here are the contents of `hip_stats.csv` file: +Here are the contents of ``hip_api_stats.csv`` file: .. csv-table:: HIP stats - :file: /data/hip_stats.csv + :file: /data/hip_api_stats.csv :widths: 10,10,20,20,10,10,10,10 :header-rows: 1 @@ -454,37 +466,37 @@ Properties .. code-block:: shell - $ cat input.txt + $ cat input.txt - pmc: GPUBusy SQ_WAVES - pmc: GRBM_GUI_ACTIVE + pmc: GPUBusy SQ_WAVES + pmc: GRBM_GUI_ACTIVE .. code-block:: shell - $ cat input.json + $ cat input.json - { - "jobs": [ - { - "pmc": ["SQ_WAVES", "GRBM_COUNT", "GUI_ACTIVE"] - }, - { - "pmc": ["FETCH_SIZE", "WRITE_SIZE"], - "kernel_include_regex": ".*_kernel", - "kernel_exclude_regex": "multiply", - "kernel_iteration_range": "[1-2]","[3-4]" - "output_file": "out", - "output_format": [ - "csv", - "json" - ], - "truncate_kernels": true - ] - } + { + "jobs": [ + { + "pmc": ["SQ_WAVES", "GRBM_COUNT", "GUI_ACTIVE"] + }, + { + "pmc": ["FETCH_SIZE", "WRITE_SIZE"], + "kernel_include_regex": ".*_kernel", + "kernel_exclude_regex": "multiply", + "kernel_iteration_range": "[1-2]","[3-4]" + "output_file": "out", + "output_format": [ + "csv", + "json" + ], + "truncate_kernels": true + ] + } .. code-block:: shell - $ cat input.yaml + $ cat input.yaml jobs: - pmc: @@ -505,19 +517,19 @@ To supply the input file for kernel profiling, use: .. code-block:: shell - rocprofv3 -i input.txt -- + rocprofv3 -i input.txt -- -Running the above command generates a `./pmc_n/counter_collection.csv` file prefixed with the process ID. For each ``pmc`` row, a directory ``pmc_n`` containing a `counter_collection.csv` file is generated, where n = 1 for the first row and so on. +Running the above command generates a ``./pmc_n/counter_collection.csv`` file prefixed with the process ID. For each ``pmc`` row, a directory ``pmc_n`` containing a ``counter_collection.csv`` file is generated, where n = 1 for the first row and so on. -In case of JSON or YAML input file, for each job, a directory ``pass_n`` containing a `counter_collection.csv` file is generated where n = 1...N jobs. +In case of JSON or YAML input file, for each job, a directory ``pass_n`` containing a ``counter_collection.csv`` file is generated where n = 1...N jobs. Each row of the CSV file is an instance of kernel execution. Here is a truncated version of the output file from ``pmc_1``: .. code-block:: shell - $ cat pmc_1/218_counter_collection.csv + $ cat pmc_1/218_counter_collection.csv -Here are the contents of `counter_collection.csv` file: +Here are the contents of ``counter_collection.csv`` file: .. csv-table:: Counter collection :file: /data/counter_collection.csv @@ -534,28 +546,28 @@ and an iteration range (set of iterations of the included kernels). If the itera .. code-block:: shell - $ cat input.yml - jobs: - - pmc: [SQ_WAVES] - kernel_include_regex: "divide" - kernel_exclude_regex: "" - kernel_iteration_range: "[1, 2, [5-8]]" + $ cat input.yml + jobs: + - pmc: [SQ_WAVES] + kernel_include_regex: "divide" + kernel_exclude_regex: "" + kernel_iteration_range: "[1, 2, [5-8]]" Agent info ++++++++++++ .. note:: - All tracing and counter collection options generate an additional `agent_info.csv` file prefixed with the process ID. + All tracing and counter collection options generate an additional ``agent_info.csv`` file prefixed with the process ID. -The `agent_info.csv` file contains information about the CPU or GPU the kernel runs on. +The ``agent_info.csv`` file contains information about the CPU or GPU the kernel runs on. .. code-block:: shell - $ cat 238_agent_info.csv + $ cat 238_agent_info.csv - "Node_Id","Logical_Node_Id","Agent_Type","Cpu_Cores_Count","Simd_Count","Cpu_Core_Id_Base","Simd_Id_Base","Max_Waves_Per_Simd","Lds_Size_In_Kb","Gds_Size_In_Kb","Num_Gws","Wave_Front_Size","Num_Xcc","Cu_Count","Array_Count","Num_Shader_Banks","Simd_Arrays_Per_Engine","Cu_Per_Simd_Array","Simd_Per_Cu","Max_Slots_Scratch_Cu","Gfx_Target_Version","Vendor_Id","Device_Id","Location_Id","Domain","Drm_Render_Minor","Num_Sdma_Engines","Num_Sdma_Xgmi_Engines","Num_Sdma_Queues_Per_Engine","Num_Cp_Queues","Max_Engine_Clk_Ccompute","Max_Engine_Clk_Fcompute","Sdma_Fw_Version","Fw_Version","Capability","Cu_Per_Engine","Max_Waves_Per_Cu","Family_Id","Workgroup_Max_Size","Grid_Max_Size","Local_Mem_Size","Hive_Id","Gpu_Id","Workgroup_Max_Dim_X","Workgroup_Max_Dim_Y","Workgroup_Max_Dim_Z","Grid_Max_Dim_X","Grid_Max_Dim_Y","Grid_Max_Dim_Z","Name","Vendor_Name","Product_Name","Model_Name" - 0,0,"CPU",24,0,0,0,0,0,0,0,0,1,24,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,3800,0,0,0,0,0,0,23,0,0,0,0,0,0,0,0,0,0,0,"AMD Ryzen 9 3900X 12-Core Processor","CPU","AMD Ryzen 9 3900X 12-Core Processor","" - 1,1,"GPU",0,256,0,2147487744,10,64,0,64,64,1,64,4,4,1,16,4,32,90000,4098,26751,12032,0,128,2,0,2,24,3800,1630,432,440,138420864,16,40,141,1024,4294967295,0,0,64700,1024,1024,1024,4294967295,4294967295,4294967295,"gfx900","AMD","Radeon RX Vega","vega10" + "Node_Id","Logical_Node_Id","Agent_Type","Cpu_Cores_Count","Simd_Count","Cpu_Core_Id_Base","Simd_Id_Base","Max_Waves_Per_Simd","Lds_Size_In_Kb","Gds_Size_In_Kb","Num_Gws","Wave_Front_Size","Num_Xcc","Cu_Count","Array_Count","Num_Shader_Banks","Simd_Arrays_Per_Engine","Cu_Per_Simd_Array","Simd_Per_Cu","Max_Slots_Scratch_Cu","Gfx_Target_Version","Vendor_Id","Device_Id","Location_Id","Domain","Drm_Render_Minor","Num_Sdma_Engines","Num_Sdma_Xgmi_Engines","Num_Sdma_Queues_Per_Engine","Num_Cp_Queues","Max_Engine_Clk_Ccompute","Max_Engine_Clk_Fcompute","Sdma_Fw_Version","Fw_Version","Capability","Cu_Per_Engine","Max_Waves_Per_Cu","Family_Id","Workgroup_Max_Size","Grid_Max_Size","Local_Mem_Size","Hive_Id","Gpu_Id","Workgroup_Max_Dim_X","Workgroup_Max_Dim_Y","Workgroup_Max_Dim_Z","Grid_Max_Dim_X","Grid_Max_Dim_Y","Grid_Max_Dim_Z","Name","Vendor_Name","Product_Name","Model_Name" + 0,0,"CPU",24,0,0,0,0,0,0,0,0,1,24,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,3800,0,0,0,0,0,0,23,0,0,0,0,0,0,0,0,0,0,0,"AMD Ryzen 9 3900X 12-Core Processor","CPU","AMD Ryzen 9 3900X 12-Core Processor","" + 1,1,"GPU",0,256,0,2147487744,10,64,0,64,64,1,64,4,4,1,16,4,32,90000,4098,26751,12032,0,128,2,0,2,24,3800,1630,432,440,138420864,16,40,141,1024,4294967295,0,0,64700,1024,1024,1024,4294967295,4294967295,4294967295,"gfx900","AMD","Radeon RX Vega","vega10" Kernel filtering +++++++++++++++++ @@ -567,24 +579,24 @@ Here is an input file with kernel filters: .. code-block:: shell - $ cat input.yml - jobs: - - pmc: [SQ_WAVES] - kernel_include_regex: "divide" - kernel_exclude_regex: "" + $ cat input.yml + jobs: + - pmc: [SQ_WAVES] + kernel_include_regex: "divide" + kernel_exclude_regex: "" To collect counters for the kernels matching the filters specified in the preceding input file, run: .. code-block:: shell - rocprofv3 -i input.yml -- + rocprofv3 -i input.yml -- - $ cat pass_1/312_counter_collection.csv - "Correlation_Id","Dispatch_Id","Agent_Id","Queue_Id","Process_Id","Thread_Id","Grid_Size","Kernel_Name","Workgroup_Size","LDS_Block_Size","Scratch_Size","VGPR_Count","SGPR_Count","Counter_Name","Counter_Value" - 4,4,1,1,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 - 8,8,1,2,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 - 12,12,1,3,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 - 16,16,1,4,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 + $ cat pass_1/312_counter_collection.csv + "Correlation_Id","Dispatch_Id","Agent_Id","Queue_Id","Process_Id","Thread_Id","Grid_Size","Kernel_Name","Workgroup_Size","LDS_Block_Size","Scratch_Size","VGPR_Count","SGPR_Count","Counter_Name","Counter_Value" + 4,4,1,1,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 + 8,8,1,2,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 + 12,12,1,3,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 + 16,16,1,4,36499,36499,1048576,"divide_kernel(float*, float const*, float const*, int, int)",64,0,0,12,16,"SQ_WAVES",16384 .. _output-file-fields: diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt index d8025a1305..a97f70e062 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt @@ -3,7 +3,7 @@ # Installation of public C++ headers (implementations) # # -set(ROCPROFILER_CXX_DETAILS_HEADER_FILES delimit.hpp join.hpp mpl.hpp name_info.hpp) +set(ROCPROFILER_CXX_DETAILS_HEADER_FILES tokenize.hpp join.hpp mpl.hpp name_info.hpp) install( FILES ${ROCPROFILER_CXX_DETAILS_HEADER_FILES} diff --git a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/delimit.hpp b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp similarity index 75% rename from projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/delimit.hpp rename to projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp index d660b515f9..46a4aa0de2 100644 --- a/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/delimit.hpp +++ b/projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp @@ -105,6 +105,64 @@ tokenize( return _result; } +/// \brief tokenize a string into a set +/// +template , + typename DelimT = std::string_view, + typename ValueT = typename ContainerT::value_type, + typename PredicateT = ValueT (*)(DelimT&&)> +inline ContainerT +tokenize( + std::string_view line, + const std::vector& delimiters, + PredicateT&& predicate = [](DelimT&& s) -> ValueT { return ValueT{s}; }) +{ + ContainerT _result = {}; + size_t _start = 0; + size_t _end = std::string::npos; + + while(_start != std::string::npos) + { + _end = std::string::npos; + + // Find the earliest occurrence of any delimiter + for(const auto& itr : delimiters) + { + size_t pos = line.find(itr, _start); + if(pos != std::string::npos && (_end == std::string::npos || pos < _end)) + { + _end = pos; + } + } + + // Extract token and update start position + if(_end != std::string::npos) + { + mpl::emplace(_result, + std::forward(predicate)(line.substr(_start, _end - _start))); + _start = _end; + + // Move start past the delimiter + for(const auto& delimiter : delimiters) + { + if(line.compare(_start, delimiter.size(), delimiter) == 0) + { + _start += delimiter.size(); + break; + } + } + } + else + { + // Last token after the final delimiter + mpl::emplace(_result, std::forward(predicate)(line.substr(_start))); + break; + } + } + + return _result; +} + /// \brief apply a string transformation to substring in between a common delimiter. /// template > diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt index 22a361dddd..cb61d3ed44 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt @@ -13,6 +13,7 @@ set(TOOL_HEADERS generateJSON.hpp generateOTF2.hpp generatePerfetto.hpp + generateStats.hpp helper.hpp output_file.hpp statistics.hpp @@ -26,9 +27,11 @@ set(TOOL_SOURCES generateJSON.cpp generateOTF2.cpp generatePerfetto.cpp + generateStats.cpp helper.cpp main.c output_file.cpp + statistics.cpp tmp_file_buffer.cpp tmp_file.cpp tool.cpp) diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp index adbac73975..cb144ad4ac 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp @@ -59,7 +59,7 @@ struct buffered_output operator bool() const { return enabled; } std::deque element_data = {}; - stats_data_t stats = {}; + stats_entry_t stats = {}; private: bool enabled = false; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp index e3744d7afc..4450d0860b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp @@ -28,12 +28,14 @@ #include "lib/common/environment.hpp" #include "lib/common/filesystem.hpp" #include "lib/common/logging.hpp" +#include "lib/common/units.hpp" #include "lib/common/utility.hpp" -#include +#include #include +#include #include #include #include @@ -44,6 +46,7 @@ #include #include #include +#include #include namespace rocprofiler @@ -67,8 +70,14 @@ const auto* launch_clock = as_pointer(std::chrono::system_clock::now()); const auto* launch_datetime = get_local_datetime(get_env("ROCP_TIME_FORMAT", "%F_%H.%M"), launch_time); const auto env_regexes = - new std::array{std::regex{"(.*)%(env|ENV)\\{([A-Z0-9_]+)\\}%(.*)"}, - std::regex{"(.*)\\$(env|ENV)\\{([A-Z0-9_]+)\\}(.*)"}}; + new std::array{std::regex{"(.*)%(env|ENV)\\{([A-Z0-9_]+)\\}%(.*)"}, + std::regex{"(.*)\\$(env|ENV)\\{([A-Z0-9_]+)\\}(.*)"}, + std::regex{"(.*)%q\\{([A-Z0-9_]+)\\}(.*)"}}; +// env regex examples: +// - %env{USER}% Consistent with other output key formats (start+end with %) +// - $ENV{USER} Similar to CMake +// - %q{USER} Compatibility with NVIDIA +// std::string* get_local_datetime(const std::string& dt_format, std::time_t*& _dt_curr) @@ -86,6 +95,22 @@ get_local_datetime(const std::string& dt_format, std::time_t*& _dt_curr) return nullptr; } +std::string +get_hostname() +{ + auto _hostname_buff = std::array{}; + _hostname_buff.fill('\0'); + if(gethostname(_hostname_buff.data(), _hostname_buff.size() - 1) != 0) + { + auto _err = errno; + ROCP_WARNING << "Hostname unknown. gethostname failed with error code " << _err << ": " + << strerror(_err); + return std::string{"UNKNOWN_HOSTNAME"}; + } + + return std::string{_hostname_buff.data()}; +} + inline bool not_is_space(int ch) { @@ -279,6 +304,34 @@ config::config() const auto supported_perfetto_backends = std::set{"inprocess", "system"}; LOG_IF(FATAL, supported_perfetto_backends.count(perfetto_backend) == 0) << "Unsupported perfetto backend type: " << perfetto_backend; + + if(stats_summary_unit == "sec") + stats_summary_unit_value = common::units::sec; + else if(stats_summary_unit == "msec") + stats_summary_unit_value = common::units::msec; + else if(stats_summary_unit == "usec") + stats_summary_unit_value = common::units::usec; + else if(stats_summary_unit == "nsec") + stats_summary_unit_value = common::units::nsec; + else + { + ROCP_FATAL << "Unsupported summary units value: " << stats_summary_unit; + } + + if(auto _summary_grps = get_env("ROCPROF_STATS_SUMMARY_GROUPS", ""); !_summary_grps.empty()) + { + stats_summary_groups = + sdk::parse::tokenize(_summary_grps, std::vector{"##@@##"}); + + // remove any empty strings (just in case these slipped through) + stats_summary_groups.erase(std::remove_if(stats_summary_groups.begin(), + stats_summary_groups.end(), + [](const auto& itr) { return itr.empty(); }), + stats_summary_groups.end()); + } + + // enable summary output if any of these are enabled + summary_output = (stats_summary || stats_summary_per_domain || !stats_summary_groups.empty()); } std::vector @@ -374,8 +427,10 @@ output_keys(std::string _tag) } auto _launch_time = (launch_datetime) ? *launch_datetime : std::string{".UNKNOWN_LAUNCH_TIME."}; + auto _hostname = get_hostname(); for(auto&& itr : std::initializer_list{ + {"%hostname%", _hostname, "Network hostname"}, {"%pid%", _proc_id, "Process identifier"}, {"%ppid%", _parent_id, "Parent process identifier"}, {"%pgid%", _pgroup_id, "Process group identifier"}, @@ -392,6 +447,7 @@ output_keys(std::string _tag) } for(auto&& itr : std::initializer_list{ + {"%h", _hostname, "Shorthand for %hostname%"}, {"%p", _proc_id, "Shorthand for %pid%"}, {"%j", _slurm_job_id, "Shorthand for %job%"}, {"%r", _slurm_proc_id, "Shorthand for %rank%"}, @@ -404,8 +460,10 @@ output_keys(std::string _tag) return _options; } +namespace +{ std::string -format(std::string _fpath, const std::string& _tag) +format_impl(std::string _fpath, const std::vector& _keys) { if(_fpath.find('%') == std::string::npos && _fpath.find('$') == std::string::npos) return _fpath; @@ -416,20 +474,36 @@ format(std::string _fpath, const std::string& _tag) _v.replace(pos, pitr.key.length(), pitr.value); }; - for(auto&& itr : output_keys(_tag)) + for(auto&& itr : _keys) _replace(_fpath, itr); // environment and configuration variables try { + auto strip_leading_and_replace = + [](std::string_view inp_v, std::initializer_list keys, const char* val) { + auto inp = std::string{inp_v}; + for(auto key : keys) + { + auto pos = std::string::npos; + while((pos = inp.find(key)) == 0) + inp = inp.substr(pos + 1); + + while((pos = inp.find(key)) != std::string::npos) + inp = inp.replace(pos, 1, val); + } + return inp; + }; + for(const auto& _re : *env_regexes) { while(std::regex_search(_fpath, _re)) { auto _var = std::regex_replace(_fpath, _re, "$3"); std::string _val = get_env(_var, ""); - auto _beg = std::regex_replace(_fpath, _re, "$1"); - auto _end = std::regex_replace(_fpath, _re, "$4"); + _val = strip_leading_and_replace(_val, {'\t', ' ', '/'}, "_"); + auto _beg = std::regex_replace(_fpath, _re, "$1"); + auto _end = std::regex_replace(_fpath, _re, "$4"); _fpath = fmt::format("{}{}{}", _beg, _val, _end); } } @@ -454,6 +528,25 @@ format(std::string _fpath, const std::string& _tag) return _fpath; } +std::string +format(std::string _fpath, const std::vector& _keys) +{ + if(_fpath.find('%') == std::string::npos && _fpath.find('$') == std::string::npos) + return _fpath; + + auto _ref = _fpath; + _fpath = format_impl(std::move(_fpath), _keys); + + return (_fpath == _ref) ? _fpath : format(std::move(_fpath), _keys); +} +} // namespace + +std::string +format(std::string _fpath, const std::string& _tag) +{ + return format(std::move(_fpath), output_keys(_tag)); +} + std::string format_name(std::string_view _name, const config& _cfg) { diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp index 8c430e9839..55f7b0d1bc 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp @@ -25,8 +25,14 @@ #include "lib/common/environment.hpp" #include "lib/common/filesystem.hpp" +#include "lib/common/mpl.hpp" + +#include + +#include #include +#include #include #include #include @@ -38,6 +44,8 @@ namespace tool namespace fs = common::filesystem; using common::get_env; +struct config; + enum class config_context { global = 0, @@ -48,6 +56,19 @@ enum class config_context perfetto_plugin, }; +void +initialize(); + +template +config& +get_config(); + +std::string +format_name(std::string_view _name, const config& = get_config<>()); + +std::string +format(std::string _fpath, const std::string& _tag = {}); + int get_mpi_size(); @@ -67,25 +88,32 @@ struct config bool hsa_finalizer_ext_api_trace = get_env("ROCPROF_HSA_FINALIZER_EXT_API_TRACE", false); bool marker_api_trace = get_env("ROCPROF_MARKER_API_TRACE", false); bool memory_copy_trace = get_env("ROCPROF_MEMORY_COPY_TRACE", false); - bool scratch_memory = get_env("ROCPROF_SCRATCH_MEMORY_TRACE", false); + bool scratch_memory_trace = get_env("ROCPROF_SCRATCH_MEMORY_TRACE", false); bool counter_collection = get_env("ROCPROF_COUNTER_COLLECTION", false); bool hip_runtime_api_trace = get_env("ROCPROF_HIP_RUNTIME_API_TRACE", false); bool hip_compiler_api_trace = get_env("ROCPROF_HIP_COMPILER_API_TRACE", false); bool list_metrics = get_env("ROCPROF_LIST_METRICS", false); bool list_metrics_output_file = get_env("ROCPROF_OUTPUT_LIST_METRICS_FILE", false); bool stats = get_env("ROCPROF_STATS", false); + bool stats_summary = get_env("ROCPROF_STATS_SUMMARY", false); + bool stats_summary_per_domain = get_env("ROCPROF_STATS_SUMMARY_PER_DOMAIN", false); bool csv_output = false; bool json_output = false; bool pftrace_output = false; bool otf2_output = false; + bool summary_output = false; bool kernel_rename = get_env("ROCPROF_KERNEL_RENAME", false); int mpi_size = get_mpi_size(); int mpi_rank = get_mpi_rank(); size_t perfetto_shmem_size_hint = get_env("ROCPROF_PERFETTO_SHMEM_SIZE_HINT_KB", 64); size_t perfetto_buffer_size = get_env("ROCPROF_PERFETTO_BUFFER_SIZE_KB", 1024000); - std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string()); - std::string output_file = get_env("ROCPROF_OUTPUT_FILE_NAME", std::to_string(getpid())); - std::string tmp_directory = get_env("ROCPROF_TMPDIR", output_path); + uint64_t stats_summary_unit_value = 1; + std::string stats_summary_unit = get_env("ROCPROF_STATS_SUMMARY_UNITS", "nsec"); + std::string output_path = get_env("ROCPROF_OUTPUT_PATH", fs::current_path().string()); + std::string output_file = + get_env("ROCPROF_OUTPUT_FILE_NAME", fmt::format("%hostname%/{}", getpid())); + std::string tmp_directory = get_env("ROCPROF_TMPDIR", output_path); + std::string stats_summary_file = get_env("ROCPROF_STATS_SUMMARY_OUTPUT", "stderr"); std::string kernel_filter_include = get_env("ROCPROF_KERNEL_FILTER_INCLUDE_REGEX", std::string{".*"}); @@ -94,11 +122,68 @@ struct config std::string perfetto_buffer_fill_policy = get_env("ROCPROF_PERFETTO_BUFFER_FILL_POLICY", std::string{"discard"}); std::string perfetto_backend = get_env("ROCPROF_PERFETTO_BACKEND", std::string{"inprocess"}); - std::unordered_set kernel_filter_range = {}; - std::set counters = {}; + std::unordered_set kernel_filter_range = {}; + std::set counters = {}; + std::vector stats_summary_groups = {}; + + template + void save(ArchiveT&) const; + + template + void load(ArchiveT&) + {} }; -template +template +void +config::save(ArchiveT& ar) const +{ +#define CFG_SERIALIZE_MEMBER(VAR) ar(cereal::make_nvp(#VAR, VAR)) +#define CFG_SERIALIZE_NAMED_MEMBER(NAME, VAR) ar(cereal::make_nvp(NAME, VAR)) + + CFG_SERIALIZE_MEMBER(demangle); + CFG_SERIALIZE_MEMBER(truncate); + CFG_SERIALIZE_MEMBER(kernel_trace); + CFG_SERIALIZE_MEMBER(hsa_core_api_trace); + CFG_SERIALIZE_MEMBER(hsa_amd_ext_api_trace); + CFG_SERIALIZE_MEMBER(hsa_image_ext_api_trace); + CFG_SERIALIZE_MEMBER(hsa_finalizer_ext_api_trace); + CFG_SERIALIZE_MEMBER(marker_api_trace); + CFG_SERIALIZE_MEMBER(memory_copy_trace); + CFG_SERIALIZE_MEMBER(scratch_memory_trace); + CFG_SERIALIZE_MEMBER(counter_collection); + CFG_SERIALIZE_MEMBER(hip_runtime_api_trace); + CFG_SERIALIZE_MEMBER(hip_compiler_api_trace); + CFG_SERIALIZE_MEMBER(kernel_rename); + + CFG_SERIALIZE_NAMED_MEMBER("summary", stats_summary); + CFG_SERIALIZE_NAMED_MEMBER("summary_per_domain", stats_summary_per_domain); + CFG_SERIALIZE_NAMED_MEMBER("summary_groups", stats_summary_groups); + CFG_SERIALIZE_NAMED_MEMBER("summary_unit", stats_summary_unit); + CFG_SERIALIZE_NAMED_MEMBER("summary_file", stats_summary_file); + + CFG_SERIALIZE_MEMBER(perfetto_shmem_size_hint); + CFG_SERIALIZE_MEMBER(perfetto_buffer_size); + CFG_SERIALIZE_MEMBER(perfetto_buffer_fill_policy); + CFG_SERIALIZE_MEMBER(perfetto_backend); + + CFG_SERIALIZE_NAMED_MEMBER("raw_tmp_directory", tmp_directory); + CFG_SERIALIZE_NAMED_MEMBER("raw_output_path", output_path); + CFG_SERIALIZE_NAMED_MEMBER("raw_output_file", output_file); + CFG_SERIALIZE_NAMED_MEMBER("tmp_directory", format(tmp_directory)); + CFG_SERIALIZE_NAMED_MEMBER("output_path", format(output_path)); + CFG_SERIALIZE_NAMED_MEMBER("output_file", format(output_file)); + + CFG_SERIALIZE_MEMBER(counters); + CFG_SERIALIZE_MEMBER(kernel_filter_include); + CFG_SERIALIZE_MEMBER(kernel_filter_exclude); + CFG_SERIALIZE_MEMBER(kernel_filter_range); + +#undef CFG_SERIALIZE_MEMBER +#undef CFG_SERIALIZE_NAMED_MEMBER +} + +template config& get_config() { @@ -119,6 +204,11 @@ struct output_key { output_key(std::string _key, std::string _val, std::string _desc = {}); + template ::value, int> = 0> + output_key(std::string _key, Tp&& _val, std::string _desc = {}); + operator std::pair() const; std::string key = {}; @@ -126,16 +216,14 @@ struct output_key std::string description = {}; }; +template ::value, int>> +output_key::output_key(std::string _key, Tp&& _val, std::string _desc) +: key{std::move(_key)} +, value{fmt::format("{}", std::forward(_val))} +, description{std::move(_desc)} +{} + std::vector output_keys(std::string _tag = {}); - -std::string -format(std::string _fpath, const std::string& _tag = {}); - -std::string -format_name(std::string_view _name, const config& = get_config<>()); - -void -initialize(); } // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp index 94750bb26f..d8cacc6f4f 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp @@ -37,13 +37,13 @@ struct domain_type_name; static constexpr auto filename = FILENAME; \ }; -DEFINE_BUFFER_TYPE_NAME(HSA, "HSA_API", "hsa_trace") -DEFINE_BUFFER_TYPE_NAME(HIP, "HIP_API", "hip_trace") -DEFINE_BUFFER_TYPE_NAME(MEMORY_COPY, "MEMORY_COPY", "memory_copy") -DEFINE_BUFFER_TYPE_NAME(COUNTER_COLLECTION, "COUNTER_COLLECTION", "counter_collection") +DEFINE_BUFFER_TYPE_NAME(HSA, "HSA_API", "hsa_api") +DEFINE_BUFFER_TYPE_NAME(HIP, "HIP_API", "hip_api") +DEFINE_BUFFER_TYPE_NAME(MARKER, "MARKER_API", "marker_api") DEFINE_BUFFER_TYPE_NAME(KERNEL_DISPATCH, "KERNEL_DISPATCH", "kernel_dispatch") -DEFINE_BUFFER_TYPE_NAME(MARKER, "MARKER", "marker_trace") +DEFINE_BUFFER_TYPE_NAME(MEMORY_COPY, "MEMORY_COPY", "memory_copy") DEFINE_BUFFER_TYPE_NAME(SCRATCH_MEMORY, "SCRATCH_MEMORY", "scratch_memory") +DEFINE_BUFFER_TYPE_NAME(COUNTER_COLLECTION, "COUNTER_COLLECTION", "counter_collection") #undef DEFINE_BUFFER_TYPE_NAME diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp index 8fa92dc403..e4b055b820 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp @@ -28,11 +28,11 @@ enum class domain_type { HSA = 0, HIP, - MEMORY_COPY, - COUNTER_COLLECTION, - KERNEL_DISPATCH, MARKER, + KERNEL_DISPATCH, + MEMORY_COPY, SCRATCH_MEMORY, + COUNTER_COLLECTION, LAST, }; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp index dd05c01115..8841a19136 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp @@ -21,9 +21,10 @@ // SOFTWARE. #include "generateCSV.hpp" +#include "config.hpp" #include "csv.hpp" +#include "generateStats.hpp" #include "helper.hpp" -#include "lib/rocprofiler-sdk-tool/config.hpp" #include "statistics.hpp" #include @@ -41,47 +42,6 @@ namespace tool { namespace { -using stats_data_t = statistics; -using stats_map_t = std::map; - -struct percentage -{ - float_type value = {}; - - friend std::ostream& operator<<(std::ostream& os, percentage val) { return (os << val.value); } -}; - -struct stats_formatter -{ - template - std::ostream& operator()(std::ostream& ofs, const Tp& _val) const - { - using value_type = common::mpl::unqualified_type_t; - - if constexpr(std::is_floating_point::value) - { - constexpr value_type one_hundredth = 1.0e-2; - if(_val > one_hundredth) - ofs << std::setprecision(6) << std::fixed; - else - ofs << std::setprecision(8) << std::scientific; - } - else if constexpr(std::is_same::value) - { - constexpr float_type one = 1.0; - constexpr float_type one_hundredth = 1.0e-2; - if(_val.value >= one) - ofs << std::setprecision(2) << std::fixed; - else if(_val.value > one_hundredth) - ofs << std::setprecision(4) << std::fixed; - else - ofs << std::setprecision(3) << std::scientific; - } - - return ofs; - } -}; - tool::output_file get_stats_output_file(std::string name) { @@ -99,10 +59,10 @@ get_stats_output_file(std::string name) }}; } -stats_data_t -write_stats(output_file&& ofs, const stats_map_t& data_v) +void +write_stats(output_file&& ofs, const stats_entry_vec_t& data_v) { - auto data = std::vector>{}; + auto data = stats_entry_vec_t{}; auto _duration = stats_data_t{}; for(const auto& [id, value] : data_v) { @@ -136,8 +96,6 @@ write_stats(output_file&& ofs, const stats_map_t& data_v) value.get_stddev()); ofs << _row.str() << std::flush; } - - return _duration; } } // namespace @@ -275,14 +233,17 @@ generate_csv(tool_table* /*tool_functions*/, std::vector } } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; - auto kernel_stats = stats_map_t{}; - auto ofs = tool::output_file{"kernel_trace", + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("kernel_stats"), stats.entries); + + auto ofs = tool::output_file{"kernel_trace", tool::csv::kernel_trace_csv_encoder{}, {"Kind", "Agent_Id", @@ -329,27 +290,21 @@ generate_csv(tool_table* record.dispatch_info.grid_size.y, record.dispatch_info.grid_size.z); - if(tool::get_config().stats) - kernel_stats[kernel_name] += (record.end_timestamp - record.start_timestamp); - ofs << row_ss.str(); } - - auto _duration = stats_data_t{}; - if(tool::get_config().stats) - _duration = write_stats(get_stats_output_file("kernel_stats"), kernel_stats); - - return _duration; } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; - auto hip_stats = stats_map_t{}; - auto ofs = tool::output_file{"hip_api_trace", + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("hip_api_stats"), stats.entries); + + auto ofs = tool::output_file{"hip_api_trace", tool::csv::api_csv_encoder{}, {"Domain", "Function", @@ -372,28 +327,21 @@ generate_csv(tool_table* tool record.start_timestamp, record.end_timestamp); - if(tool::get_config().stats) - hip_stats[api_name] += (record.end_timestamp - record.start_timestamp); - ofs << row_ss.str(); } - - auto _duration = stats_data_t{}; - if(tool::get_config().stats) - { - _duration = write_stats(get_stats_output_file("hip_stats"), hip_stats); - } - return _duration; } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; - auto hsa_stats = stats_map_t{}; - auto ofs = tool::output_file{"hsa_api_trace", + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("hsa_api_stats"), stats.entries); + + auto ofs = tool::output_file{"hsa_api_trace", tool::csv::api_csv_encoder{}, {"Domain", "Function", @@ -417,28 +365,21 @@ generate_csv(tool_table* tool record.start_timestamp, record.end_timestamp); - if(tool::get_config().stats) - hsa_stats[api_name] += (record.end_timestamp - record.start_timestamp); - ofs << row_ss.str(); } - - auto _duration = stats_data_t{}; - if(tool::get_config().stats) - { - _duration = write_stats(get_stats_output_file("hsa_stats"), hsa_stats); - } - return _duration; } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; - auto memory_copy_stats = stats_map_t{}; - auto ofs = tool::output_file{"memory_copy_trace", + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("memory_copy_stats"), stats.entries); + + auto ofs = tool::output_file{"memory_copy_trace", tool::csv::memory_copy_csv_encoder{}, {"Kind", "Direction", @@ -461,28 +402,21 @@ generate_csv(tool_table* record.start_timestamp, record.end_timestamp); - if(tool::get_config().stats) - memory_copy_stats[api_name] += (record.end_timestamp - record.start_timestamp); - ofs << row_ss.str(); } - - auto _duration = stats_data_t{}; - if(tool::get_config().stats) - { - _duration = write_stats(get_stats_output_file("memory_copy_stats"), memory_copy_stats); - } - return _duration; } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; - auto marker_stats = stats_map_t{}; - auto ofs = tool::output_file{"marker_api_trace", + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("marker_api_stats"), stats.entries); + + auto ofs = tool::output_file{"marker_api_trace", tool::csv::marker_csv_encoder{}, {"Domain", "Function", @@ -518,25 +452,19 @@ generate_csv(tool_table* t record.start_timestamp, record.end_timestamp); - if(tool::get_config().stats) - marker_stats[_name] += (record.end_timestamp - record.start_timestamp); - ofs << row_ss.str(); } - - auto _duration = stats_data_t{}; - if(tool::get_config().stats) - { - _duration = write_stats(get_stats_output_file("marker_stats"), marker_stats); - } - return _duration; } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; + + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("counter_collection_stats"), stats.entries); auto ofs = tool::output_file{"counter_collection", tool::csv::counter_collection_csv_encoder{}, @@ -601,14 +529,17 @@ generate_csv(tool_table* too } ofs << row_ss.str(); } - return stats_data_t{}; } -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data) + const std::deque& data, + const stats_entry_t& stats) { - if(data.empty()) return stats_data_t{}; + if(data.empty()) return; + + if(tool::get_config().stats && stats) + write_stats(get_stats_output_file("scratch_memory_stats"), stats.entries); auto ofs = tool::output_file{"scratch_memory_trace", tool::csv::scratch_memory_encoder{}, @@ -623,7 +554,6 @@ generate_csv(tool_table* "End_Timestamp", }}; - auto scratch_memory_stats = stats_map_t{}; for(const auto& record : data) { auto row_ss = std::stringstream{}; @@ -641,45 +571,38 @@ generate_csv(tool_table* record.start_timestamp, record.end_timestamp); - if(tool::get_config().stats) - scratch_memory_stats[op_name] += (record.end_timestamp - record.start_timestamp); - ofs << row_ss.str(); } - - auto _duration = stats_data_t{}; - if(tool::get_config().stats) - { - _duration = - write_stats(get_stats_output_file("scratch_memory_stats"), scratch_memory_stats); - } - return _duration; } void -generate_csv(tool_table* /*tool_functions*/, std::unordered_map& data) +generate_csv(tool_table* /*tool_functions*/, const domain_stats_vec_t& data_v) { using csv_encoder_t = rocprofiler::tool::csv::stats_csv_encoder; if(!tool::get_config().stats) return; + auto _data = data_v; auto _total_stats = stats_data_t{}; - for(const auto& itr : data) - _total_stats += itr.second; + for(const auto& itr : _data) + _total_stats += itr.second.total; if(_total_stats.get_count() == 0) return; + std::sort(_data.begin(), _data.end(), [](const auto& lhs, const auto& rhs) { + return (lhs.second.total.get_sum() > rhs.second.total.get_sum()); + }); + auto ofs = get_stats_output_file("domain_stats"); - constexpr float_type one_hundred = 100.0; - const float_type _total_duration = _total_stats.get_sum(); - for(const auto& [type, value] : data) + const float_type _total_duration = _total_stats.get_sum(); + for(const auto& [type, value] : _data) { - auto name = get_domain_column_name(type); - auto duration_ns = value.get_sum(); - auto calls = value.get_count(); - float_type avg_ns = value.get_mean(); - float_type percent_v = (duration_ns / _total_duration) * one_hundred; + auto name = get_domain_column_name(type); + auto duration_ns = value.total.get_sum(); + auto calls = value.total.get_count(); + auto avg_ns = value.total.get_mean(); + auto percent_v = value.total.get_percent(_total_duration); auto _row = std::stringstream{}; csv_encoder_t::write_row(_row, @@ -688,9 +611,9 @@ generate_csv(tool_table* /*tool_functions*/, std::unordered_map; - void generate_csv(tool_table* tool_functions, std::vector& data); -stats_data_t +void generate_csv(tool_table* tool_functions, - const std::deque& data); - -stats_data_t -generate_csv(tool_table* tool_functions, - const std::deque& data); - -stats_data_t -generate_csv(tool_table* tool_functions, - const std::deque& data); - -stats_data_t -generate_csv(tool_table* tool_functions, - const std::deque& data); - -stats_data_t -generate_csv(tool_table* tool_functions, - const std::deque& data); - -stats_data_t -generate_csv(tool_table* tool_functions, - const std::deque& data); - -stats_data_t -generate_csv(tool_table* tool_functions, - const std::deque& data); + const std::deque& data, + const stats_entry_t& stats); void -generate_csv(tool_table* tool_functions, std::unordered_map& data); +generate_csv(tool_table* tool_functions, + const std::deque& data, + const stats_entry_t& stats); + +void +generate_csv(tool_table* tool_functions, + const std::deque& data, + const stats_entry_t& stats); + +void +generate_csv(tool_table* tool_functions, + const std::deque& data, + const stats_entry_t& stats); + +void +generate_csv(tool_table* tool_functions, + const std::deque& data, + const stats_entry_t& stats); + +void +generate_csv(tool_table* tool_functions, + const std::deque& data, + const stats_entry_t& stats); + +void +generate_csv(tool_table* tool_functions, + const std::deque& data, + const stats_entry_t& stats); + +void +generate_csv(tool_table* tool_functions, const domain_stats_vec_t& data); } // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp index b0ecec970f..2256d826c6 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp @@ -24,8 +24,10 @@ #include "config.hpp" #include "helper.hpp" #include "output_file.hpp" +#include "statistics.hpp" #include "lib/common/string_entry.hpp" +#include "lib/common/utility.hpp" #include #include @@ -39,6 +41,7 @@ namespace tool void write_json(tool_table* tool_functions, uint64_t pid, + const domain_stats_vec_t& domain_stats, std::vector agent_data, std::vector counter_data, std::deque* hip_api_deque, @@ -52,14 +55,14 @@ write_json(tool_table* tool { using JSONOutputArchive = cereal::MinimalJSONOutputArchive; - constexpr auto json_prec = 32; - constexpr auto json_indent = JSONOutputArchive::Options::IndentChar::space; - auto json_opts = JSONOutputArchive::Options{json_prec, json_indent, 1}; - auto filename = std::string_view{"results"}; - auto [output_stream, cleanup] = get_output_stream(filename, ".json"); + constexpr auto json_prec = 32; + constexpr auto json_indent = JSONOutputArchive::Options::IndentChar::space; + auto json_opts = JSONOutputArchive::Options{json_prec, json_indent, 1}; + auto filename = std::string_view{"results"}; + auto ofs = get_output_stream(filename, ".json"); { - auto json_ar = JSONOutputArchive{*output_stream, json_opts}; + auto json_ar = JSONOutputArchive{*ofs.stream, json_opts}; json_ar.setNextName("rocprofiler-sdk-tool"); json_ar.startNode(); @@ -74,6 +77,29 @@ write_json(tool_table* tool json_ar(cereal::make_nvp("pid", pid)); json_ar(cereal::make_nvp("init_time", timestamps->app_start_time)); json_ar(cereal::make_nvp("fini_time", timestamps->app_end_time)); + json_ar(cereal::make_nvp("config", get_config())); + json_ar(cereal::make_nvp("command", common::read_command_line(getpid()))); + json_ar.finishNode(); + } + + // summary + { + json_ar.setNextName("summary"); + json_ar.startNode(); + json_ar.makeArray(); + + for(const auto& itr : domain_stats) + { + auto _name = get_domain_column_name(itr.first); + json_ar.startNode(); + + json_ar(cereal::make_nvp("domain", std::string{_name})); + json_ar(cereal::make_nvp("stats", itr.second)); + // itr.second.serialize(json_ar, 0); + + json_ar.finishNode(); + } + json_ar.finishNode(); } @@ -154,8 +180,7 @@ write_json(tool_table* tool json_ar.finishNode(); } - *output_stream << std::flush; - if(cleanup) cleanup(output_stream); + ofs.close(); } } // namespace tool diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp index bb40cc17c8..60ea730472 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp @@ -23,6 +23,7 @@ #pragma once #include "helper.hpp" +#include "statistics.hpp" namespace rocprofiler { @@ -31,6 +32,7 @@ namespace tool void write_json(tool_table* tool_functions, uint64_t pid, + const domain_stats_vec_t& domain_stats, std::vector agent_data, std::vector counter_data, std::deque* hip_api_deque, diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateOTF2.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateOTF2.cpp index e82903ee9c..f615d70f52 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateOTF2.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateOTF2.cpp @@ -21,13 +21,14 @@ // SOFTWARE. #include "generateOTF2.hpp" +#include "config.hpp" #include "helper.hpp" +#include "output_file.hpp" + #include "lib/common/filesystem.hpp" #include "lib/common/mpl.hpp" #include "lib/common/units.hpp" #include "lib/common/utility.hpp" -#include "lib/rocprofiler-sdk-tool/config.hpp" -#include "output_file.hpp" #include #include @@ -510,6 +511,10 @@ write_otf2(tool_table* tool if(!_inp) return; for(auto itr : *_inp) { + if(itr.kind == ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API && + itr.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA) + continue; + using value_type = common::mpl::unqualified_type_t; auto name = buffer_names.at(itr.kind, itr.operation); auto paradigm = OTF2_PARADIGM_HIP; diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp index 638faed923..f7962bdf72 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp @@ -21,11 +21,11 @@ // SOFTWARE. #include "generatePerfetto.hpp" +#include "config.hpp" #include "helper.hpp" #include "output_file.hpp" #include "lib/common/utility.hpp" -#include "lib/rocprofiler-sdk-tool/config.hpp" #include #include @@ -503,23 +503,20 @@ write_perfetto( tracing_session->FlushBlocking(); tracing_session->StopBlocking(); - auto filename = std::string{"results"}; - auto cleanup = std::function{}; - std::ostream* ofs = nullptr; - - std::tie(ofs, cleanup) = get_output_stream(filename, ".pftrace"); + auto filename = std::string{"results"}; + auto ofs = get_output_stream(filename, ".pftrace"); auto amount_read = std::atomic{0}; auto is_done = std::promise{}; auto _mtx = std::mutex{}; - auto _reader = [ofs, &_mtx, &is_done, &amount_read]( + auto _reader = [&ofs, &_mtx, &is_done, &amount_read]( ::perfetto::TracingSession::ReadTraceCallbackArgs _args) { auto _lk = std::unique_lock{_mtx}; if(_args.data && _args.size > 0) { ROCP_TRACE << "Writing " << _args.size << " B to trace..."; // Write the trace data into file - ofs->write(_args.data, _args.size); + ofs.stream->write(_args.data, _args.size); amount_read += _args.size; } ROCP_INFO_IF(!_args.has_more && amount_read > 0) @@ -540,10 +537,10 @@ write_perfetto( tracing_session.reset(); ROCP_TRACE << "Flushing trace output stream..."; - (*ofs) << std::flush; + (*ofs.stream) << std::flush; ROCP_TRACE << "Destroying trace output stream..."; - if(cleanup) cleanup(ofs); + ofs.close(); } } // namespace tool diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.cpp new file mode 100644 index 0000000000..031ca17912 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.cpp @@ -0,0 +1,350 @@ +// 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 "generateStats.hpp" +#include "config.hpp" +#include "helper.hpp" +#include "lib/common/logging.hpp" +#include "lib/rocprofiler-sdk-tool/domain_type.hpp" +#include "output_file.hpp" +#include "statistics.hpp" + +#include +#include +#include + +#include +#include +#include +#include +#include +#include +#include + +namespace rocprofiler +{ +namespace tool +{ +namespace +{ +stats_entry_t +get_stats(const stats_map_t& data_v) +{ + auto _stats = stats_entry_t{}; + for(const auto& [id, value] : data_v) + { + _stats.entries.emplace_back(id, value); + _stats.total += value; + } + + return _stats.sort(); +} +} // namespace + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data) +{ + auto kernel_stats = stats_map_t{}; + for(const auto& record : data) + { + auto kernel_name = tool_functions->tool_get_kernel_name_fn( + record.dispatch_info.kernel_id, record.correlation_id.external.value); + + kernel_stats[kernel_name] += (record.end_timestamp - record.start_timestamp); + } + + return get_stats(kernel_stats); +} + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data) +{ + auto hip_stats = stats_map_t{}; + for(const auto& record : data) + { + auto api_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + hip_stats[api_name] += (record.end_timestamp - record.start_timestamp); + } + + return get_stats(hip_stats); +} + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data) +{ + auto hsa_stats = stats_map_t{}; + for(const auto& record : data) + { + auto api_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + hsa_stats[api_name] += (record.end_timestamp - record.start_timestamp); + } + + return get_stats(hsa_stats); +} + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data) +{ + auto memory_copy_stats = stats_map_t{}; + for(const auto& record : data) + { + auto api_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + memory_copy_stats[api_name] += (record.end_timestamp - record.start_timestamp); + } + + return get_stats(memory_copy_stats); +} + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data) +{ + auto marker_stats = stats_map_t{}; + for(const auto& record : data) + { + auto _name = std::string_view{}; + + if(record.kind == ROCPROFILER_BUFFER_TRACING_MARKER_CORE_API && + (record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxMarkA || + record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangePushA || + record.operation == ROCPROFILER_MARKER_CORE_API_ID_roctxRangeStartA)) + { + _name = tool_functions->tool_get_roctx_msg_fn(record.correlation_id.internal); + } + else + { + _name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + } + + marker_stats[_name] += (record.end_timestamp - record.start_timestamp); + } + + return get_stats(marker_stats); +} + +stats_entry_t +generate_stats(tool_table* /*tool_functions*/, + const std::deque& /*data*/) +{ + return stats_entry_t{}; +} + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data) +{ + auto scratch_memory_stats = stats_map_t{}; + for(const auto& record : data) + { + auto op_name = tool_functions->tool_get_operation_name_fn(record.kind, record.operation); + scratch_memory_stats[op_name] += (record.end_timestamp - record.start_timestamp); + } + + return get_stats(scratch_memory_stats); +} + +namespace +{ +void +generate_stats(output_stream_t& os, + std::string_view label, + const domain_stats_vec_t& data_v, + std::string_view indent_v) +{ + auto _data = stats_entry_t{}; + auto _cols = std::unordered_map{}; + + auto _get_entry = [&_data, &_cols](domain_type _domain, + std::string_view _key) -> stats_data_t* { + for(auto& itr : _data.entries) + { + if(itr.first == _key) return &itr.second; + } + + _cols.emplace(_key, _domain); + auto& itr = _data.entries.emplace_back(_key, stats_data_t{}); + return &itr.second; + }; + + uint64_t name_width = 40; + uint64_t domain_width = 12; + for(const auto& itr : data_v) + { + for(const auto& eitr : itr.second.entries) + { + _data.total += eitr.second; + auto* _entry = _get_entry(itr.first, eitr.first); + *CHECK_NOTNULL(_entry) += eitr.second; + name_width = std::max(name_width, eitr.first.length()); + } + domain_width = std::max(domain_width, get_domain_column_name(itr.first).length()); + } + + if(!_data) return; + + std::sort(_data.entries.begin(), _data.entries.end(), [](const auto& lhs, const auto& rhs) { + return (lhs.second.get_sum() > rhs.second.get_sum()); + }); + + const float_type _total_duration = _data.total.get_sum(); + + os << fmt::format("\n{}ROCPROFV3 {}:\n\n", indent_v, label) << std::flush; + + { + auto _header = fmt::format( + "| {:^{}} | {:^{}} | {:^15} | {:^15} | {:^15} | {:^13} | {:^15} | {:^15} | {:^15} |", + "NAME", + name_width, + "DOMAIN", + domain_width, + "CALLS", + fmt::format("DURATION ({})", tool::get_config().stats_summary_unit), + fmt::format("AVERAGE ({})", tool::get_config().stats_summary_unit), + "PERCENT (INC)", + fmt::format("MIN ({})", tool::get_config().stats_summary_unit), + fmt::format("MAX ({})", tool::get_config().stats_summary_unit), + "STDDEV"); + (*os.stream) << indent_v << _header << "\n" << std::flush; + + auto _div = + fmt::format("|-{0:-^{1}}-|-{0:-^{2}}-|-{0:-^15}-|-{0:-^15}-|-{0:-^15}-|-{0:-^13}" + "-|-{0:-^15}-|-{0:-^15}-|-{0:-^15}-|", + "", + name_width, + domain_width); + (*os.stream) << indent_v << _div << "\n" << std::flush; + } + + for(const auto& [type, value] : _data.entries) + { + auto name = type; + auto duration_ns = value.get_sum(); + auto calls = value.get_count(); + auto avg_ns = value.get_mean(); + auto percent_v = value.get_percent(_total_duration); + auto percent = std::to_string(percent_v); + + auto _row = std::string{}; + + if(tool::get_config().stats_summary_unit_value > 1) + { + auto _unit_div = static_cast(tool::get_config().stats_summary_unit_value); + _row = fmt::format("{}| {:<{}} | {:<{}} | {:15} | {:15} | {:15.3e} | {:>13} | {:15} | " + "{:15} | {:15.3e} |", + indent_v, + name, + name_width, + get_domain_column_name(_cols.at(name)), + domain_width, + calls, + duration_ns / _unit_div, + avg_ns / _unit_div, + percent, + value.get_min() / _unit_div, + value.get_max() / _unit_div, + value.get_stddev() / _unit_div); + } + else + { + _row = fmt::format("{}| {:<{}} | {:<{}} | {:15} | {:15} | {:15.3e} | {:>13} | {:15} | " + "{:15} | {:15.3e} |", + indent_v, + name, + name_width, + get_domain_column_name(_cols.at(name)), + domain_width, + calls, + duration_ns, + avg_ns, + percent, + value.get_min(), + value.get_max(), + value.get_stddev()); + } + + (*os.stream) << _row << "\n" << std::flush; + } + + (*os.stream) << "\n" << std::flush; +} +} // namespace + +void +generate_stats(tool_table* /*tool_functions*/, const domain_stats_vec_t& inp_data) +{ + auto data_v = inp_data; + + std::sort(data_v.begin(), data_v.end(), [](const auto& lhs, const auto& rhs) { + return lhs.first < rhs.first; + }); + + output_stream_t _os = get_output_stream(tool::get_config().stats_summary_file, ".txt"); + auto _indent = (_os.writes_to_file()) ? std::string_view{} : std::string_view{" "}; + + if(tool::get_config().stats_summary_per_domain) + { + for(const auto& itr : data_v) + { + if(!itr.second) continue; + + auto _name = fmt::format("{} SUMMARY", get_domain_column_name(itr.first)); + auto _tmp = domain_stats_vec_t{}; + _tmp.emplace_back(itr.first, itr.second); + generate_stats(_os, _name, _tmp, _indent); + } + } + + if(!tool::get_config().stats_summary_groups.empty()) + { + auto domain_groups = std::vector{}; + for(const auto& itr : tool::get_config().stats_summary_groups) + { + auto _names = std::vector{}; + auto _tmp = domain_stats_vec_t{}; + for(const auto& ditr : data_v) + { + auto _col_name = get_domain_column_name(ditr.first); + + if(std::regex_match(_col_name.data(), std::regex{itr})) + { + if(!ditr.second) continue; + _names.emplace_back(_col_name); + _tmp.emplace_back(ditr.first, ditr.second); + } + } + + ROCP_CI_LOG_IF(WARNING, _names.empty()) + << "summary group regex '" << itr << "' matched with zero domain groups"; + + auto _name = fmt::format("{} SUMMARY", fmt::join(_names.begin(), _names.end(), " + ")); + generate_stats(_os, _name, _tmp, _indent); + } + } + + if(tool::get_config().stats_summary) generate_stats(_os, "SUMMARY", data_v, _indent); +} +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.hpp new file mode 100644 index 0000000000..f9124bdb59 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.hpp @@ -0,0 +1,63 @@ +// 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 "helper.hpp" +#include "statistics.hpp" + +namespace rocprofiler +{ +namespace tool +{ +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +stats_entry_t +generate_stats(tool_table* tool_functions, + const std::deque& data); + +void +generate_stats(tool_table* tool_functions, const domain_stats_vec_t& data); +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp index 8d04e7017b..9b76ea70d4 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp @@ -33,6 +33,12 @@ namespace tool { namespace fs = common::filesystem; +namespace +{ +const auto stdout_names = std::unordered_set{"stdout", "STDOUT"}; +const auto stderr_names = std::unordered_set{"stderr", "STDERR"}; +} // namespace + std::string get_output_filename(std::string_view fname, std::string_view ext) { @@ -55,16 +61,17 @@ get_output_filename(std::string_view fname, std::string_view ext) return tool::format(output_path / fmt::format("{}_{}{}", output_prefix, fname, _ext)); } -std::pair + +output_stream_t get_output_stream(std::string_view fname, std::string_view ext) { auto cfg_output_path = tool::format(tool::get_config().output_path); - if(cfg_output_path == "stdout" || cfg_output_path == "STDOUT") + if(stdout_names.count(cfg_output_path) > 0 || stdout_names.count(fname) > 0) return {&std::cout, [](auto*&) {}}; - else if(cfg_output_path == "stderr" || cfg_output_path == "STDERR") + else if(stderr_names.count(cfg_output_path) > 0 || stderr_names.count(fname) > 0) return {&std::cout, [](auto*&) {}}; - else if(cfg_output_path.empty()) + else if(cfg_output_path.empty() || fname.empty()) return {&std::clog, [](auto*&) {}}; auto output_file = get_output_filename(fname, ext); @@ -82,12 +89,12 @@ get_output_stream(std::string_view fname, std::string_view ext) output_file::~output_file() { - if(m_stream) + if(m_os.stream) ROCP_INFO << "Closing result file: " << m_name; else ROCP_WARNING << "output_file::~output_file does not have a output stream instance!"; - m_dtor(m_stream); + m_os.close(); } } // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp index 40e08a9c8c..c2f7ce273b 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp @@ -26,7 +26,6 @@ #include "csv.hpp" #include "lib/common/filesystem.hpp" -#include "lib/rocprofiler-sdk-tool/csv.hpp" #include #include @@ -41,12 +40,48 @@ namespace rocprofiler { namespace tool { -using output_stream_dtor_t = void (*)(std::ostream*&); +using ostream_dtor_t = void (*)(std::ostream*&); + +using output_stream_pair_t = std::pair; + +struct output_stream_t +{ + output_stream_t() = default; + output_stream_t(std::ostream* _os, ostream_dtor_t _dtor) + : stream{_os} + , dtor{_dtor} + {} + + ~output_stream_t() { close(); } + output_stream_t(const output_stream_t&) = delete; + output_stream_t(output_stream_t&&) noexcept = default; + output_stream_t& operator=(const output_stream_t&) = delete; + output_stream_t& operator=(output_stream_t&&) noexcept = default; + + explicit operator bool() const { return stream != nullptr; } + + template + std::ostream& operator<<(Tp&& value) + { + return ((stream) ? *stream : std::cerr) << std::forward(value) << std::flush; + } + + void close() + { + if(stream) (*stream) << std::flush; + if(dtor) dtor(stream); + } + + bool writes_to_file() const { return (dynamic_cast(stream) != nullptr); } + + std::ostream* stream = nullptr; + ostream_dtor_t dtor = nullptr; +}; std::string get_output_filename(std::string_view fname, std::string_view ext); -std::pair +output_stream_t get_output_stream(std::string_view fname, std::string_view ext); struct output_file @@ -65,16 +100,15 @@ struct output_file std::ostream& operator<<(T&& value) { auto _lk = std::unique_lock{m_mutex}; - return ((m_stream) ? *m_stream : std::cerr) << std::forward(value) << std::flush; + return ((m_os.stream) ? *m_os.stream : std::cerr) << std::forward(value) << std::flush; } - operator bool() const { return m_stream != nullptr; } + operator bool() const { return m_os.stream != nullptr; } private: - const std::string m_name = {}; - std::mutex m_mutex = {}; - std::ostream* m_stream = nullptr; - output_stream_dtor_t m_dtor = [](std::ostream*&) {}; + const std::string m_name = {}; + std::mutex m_mutex = {}; + output_stream_t m_os = {}; }; template @@ -82,9 +116,8 @@ output_file::output_file(std::string name, csv::csv_encoder encoder, std::array&& header) : m_name{std::move(name)} +, m_os{get_output_stream(m_name, ".csv")} { - std::tie(m_stream, m_dtor) = get_output_stream(m_name, ".csv"); - for(auto& itr : header) { ROCP_FATAL_IF(itr.empty()) @@ -92,7 +125,7 @@ output_file::output_file(std::string name, } // write the csv header - if(m_stream) encoder.write_row(*m_stream, header); + if(m_os.stream) encoder.write_row(*m_os.stream, header); } } // namespace tool } // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.cpp new file mode 100644 index 0000000000..e04e9837f8 --- /dev/null +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.cpp @@ -0,0 +1,75 @@ +// 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 "statistics.hpp" + +namespace rocprofiler +{ +namespace tool +{ +template +std::ostream& +stats_formatter::operator()(std::ostream& ofs, const Tp& _val) const +{ + using value_type = common::mpl::unqualified_type_t; + + if constexpr(std::is_floating_point::value) + { + constexpr value_type one_hundredth = 1.0e-2; + if(_val > one_hundredth) + ofs << std::setprecision(6) << std::fixed; + else + ofs << std::setprecision(8) << std::scientific; + } + else if constexpr(std::is_same::value) + { + constexpr float_type one = 1.0; + constexpr float_type one_hundredth = 1.0e-2; + if(_val.value >= one) + ofs << std::setprecision(2) << std::fixed; + else if(_val.value > one_hundredth) + ofs << std::setprecision(4) << std::fixed; + else + ofs << std::setprecision(3) << std::scientific; + } + + return ofs; +} + +#define STATS_FORMATTER_INSTANTIATE_TEMPLATE(TYPE) \ + template std::ostream& stats_formatter::operator()(std::ostream&, const TYPE&) const; + +STATS_FORMATTER_INSTANTIATE_TEMPLATE(std::string) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(std::string_view) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(percentage) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(float) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(double) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(uint8_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(uint16_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(uint32_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(uint64_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(int8_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(int16_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(int32_t) +STATS_FORMATTER_INSTANTIATE_TEMPLATE(int64_t) +} // namespace tool +} // namespace rocprofiler diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.hpp index 1913ebde1e..8e2694f8d9 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.hpp @@ -22,6 +22,13 @@ #pragma once +#include "domain_type.hpp" + +#include "lib/common/logging.hpp" +#include "lib/common/mpl.hpp" + +#include + #include #include #include @@ -29,7 +36,9 @@ #include #include #include +#include #include +#include namespace rocprofiler { @@ -93,6 +102,8 @@ public: } float_type get_stddev() const { return ::std::sqrt(::std::abs(get_variance())); } + float_type get_percent(float_type _total) const; + float_type get_percent(const this_type&) const; // Modifications void reset() @@ -209,6 +220,107 @@ public: { return statistics(lhs) -= rhs; } + + template + void serialize(ArchiveT& ar, const unsigned int) const + { + ar(cereal::make_nvp("count", m_cnt)); + ar(cereal::make_nvp("sum", m_sum)); + ar(cereal::make_nvp("sqr", m_sqr)); + ar(cereal::make_nvp("min", m_min)); + ar(cereal::make_nvp("max", m_max)); + ar(cereal::make_nvp("mean", get_mean())); + ar(cereal::make_nvp("stddev", get_stddev())); + ar(cereal::make_nvp("variance", get_variance())); + } +}; + +template +typename statistics::float_type +statistics::get_percent(float_type _total) const +{ + constexpr float_type one_hundred = 100.0; + const float_type _sum = get_sum(); + + ROCP_WARNING_IF(static_cast(_sum) > static_cast(_total)) + << "percentage calculation > 100%. sum=" << _sum << " > total=" << _total; + + return (_sum / _total) * one_hundred; +} + +template +typename statistics::float_type +statistics::get_percent(const statistics& _rhs) const +{ + return get_percent(_rhs.get_sum()); +} + +using float_type = double; +using stats_data_t = statistics; +using stats_map_t = std::map; +using stats_pair_t = std::pair; +using stats_entry_vec_t = std::vector; + +inline bool +default_stats_sorter(const stats_pair_t& lhs, const stats_pair_t& rhs) +{ + return (lhs.second.get_sum() > rhs.second.get_sum()); +} + +struct stats_entry_t +{ + using sort_predicate_t = bool (*)(const stats_pair_t&, const stats_pair_t&); + + stats_entry_t() = default; + ~stats_entry_t() = default; + stats_entry_t(const stats_entry_t&) = default; + stats_entry_t(stats_entry_t&&) noexcept = default; + stats_entry_t& operator=(const stats_entry_t&) = default; + stats_entry_t& operator=(stats_entry_t&&) noexcept = default; + + template + stats_entry_t& sort(FuncT&& _predicate = default_stats_sorter); + + explicit operator bool() const { return (total.get_count() > 0 && !entries.empty()); } + + stats_data_t total = {}; + stats_entry_vec_t entries = {}; + + template + void serialize(ArchiveT& ar, const unsigned int) const + { + total.serialize(ar, 0); + auto _entries_map = std::map{}; + for(const auto& itr : entries) + _entries_map.emplace(std::string{itr.first}, itr.second); + ar(cereal::make_nvp("operations", _entries_map)); + } +}; + +template +stats_entry_t& +stats_entry_t::sort(FuncT&& _predicate) +{ + std::sort(entries.begin(), entries.end(), std::forward(_predicate)); + return *this; +} + +using domain_stats_t = std::pair; +using domain_stats_vec_t = std::vector; + +struct stats_formatter +{ + template + std::ostream& operator()(std::ostream& ofs, const Tp& _val) const; +}; + +struct percentage +{ + float_type value = {}; + friend std::ostream& operator<<(std::ostream& os, percentage val) + { + return (stats_formatter{}(os, val) << val.value); + } }; } // namespace tool } // namespace rocprofiler @@ -228,4 +340,12 @@ min(::rocprofiler::tool::statistics lhs, const Tp& rhs) { return lhs.get_min(rhs); } + +inline std::string +to_string(::rocprofiler::tool::percentage val) +{ + auto _ss = std::stringstream{}; + _ss << val; + return _ss.str(); +} } // namespace std diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp index 67152f7fa6..ca80e569d6 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp @@ -22,6 +22,7 @@ #pragma once +#include "domain_type.hpp" #include "helper.hpp" #include "tmp_file.hpp" @@ -35,6 +36,7 @@ #include #include #include +#include #include template @@ -71,15 +73,60 @@ void write_ring_buffer(Tp _v, domain_type type) { auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer>(type); - if(_tmp_buf->capacity() == 0) return; + + if(_tmp_buf->capacity() == 0) + { + ROCP_INFO << "rocprofv3 is dropping record from domain " << get_domain_column_name(type) + << ". Buffer has a capacity of zero."; + return; + } + auto* ptr = _tmp_buf->request(false); if(ptr == nullptr) { offload_buffer>(type); ptr = _tmp_buf->request(false); - CHECK(ptr != nullptr); + + // if failed, try again + if(!ptr) ptr = _tmp_buf->request(false); + + // after second failure, emit warning message + ROCP_CI_LOG_IF(WARNING, !ptr) + << "rocprofv3 is dropping record from domain " << get_domain_column_name(type) + << ". No space in buffer: " + << fmt::format( + "capacity={}, record_size={}, used_count={}, free_count={} | raw_info=[{}]", + _tmp_buf->capacity(), + _tmp_buf->data_size(), + _tmp_buf->count(), + _tmp_buf->free(), + _tmp_buf->as_string()); + } + + if(ptr) + { + if constexpr(std::is_move_constructible::value) + { + new(ptr) Tp{std::move(_v)}; + } + else if constexpr(std::is_move_assignable::value) + { + *ptr = std::move(_v); + } + else if constexpr(std::is_copy_constructible::value) + { + new(ptr) Tp{_v}; + } + else if constexpr(std::is_copy_assignable::value) + { + *ptr = _v; + } + else + { + static_assert(std::is_void::value, + "data type is neither move/copy constructible nor move/copy assignable"); + } } - *ptr = std::move(_v); } template diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp index 47ada1485e..54f0342132 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp @@ -28,8 +28,10 @@ #include "generateJSON.hpp" #include "generateOTF2.hpp" #include "generatePerfetto.hpp" +#include "generateStats.hpp" #include "helper.hpp" #include "output_file.hpp" +#include "statistics.hpp" #include "tmp_file.hpp" #include "lib/common/environment.hpp" @@ -1335,7 +1337,7 @@ tool_init(rocprofiler_client_finalize_t fini_func, void* tool_data) "buffer tracing service for memory copy configure"); } - if(tool::get_config().scratch_memory) + if(tool::get_config().scratch_memory_trace) { ROCPROFILER_CALL(rocprofiler_create_buffer(get_client_ctx(), buffer_size, @@ -1500,21 +1502,32 @@ api_registration_callback(rocprofiler_intercept_table_t, "Iterate rocporfiler agents") } -using stats_data_t = ::rocprofiler::tool::stats_data_t; +using stats_data_t = ::rocprofiler::tool::stats_data_t; +using stats_entry_t = ::rocprofiler::tool::stats_entry_t; +using domain_stats_vec_t = ::rocprofiler::tool::domain_stats_vec_t; template void generate_output(rocprofiler::tool::buffered_output& output_v, - std::unordered_map& contributions_v) + domain_stats_vec_t& contributions_v) { if(!output_v) return; output_v.read(); + if(tool::get_config().stats || tool::get_config().summary_output) + { + output_v.stats = rocprofiler::tool::generate_stats(tool_functions, output_v.element_data); + } + + if(output_v.stats) + { + contributions_v.emplace_back(output_v.buffer_type_v, output_v.stats); + } + if(tool::get_config().csv_output) { - output_v.stats = rocprofiler::tool::generate_csv(tool_functions, output_v.element_data); - contributions_v.emplace(output_v.buffer_type_v, output_v.stats); + rocprofiler::tool::generate_csv(tool_functions, output_v.element_data, output_v.stats); } } @@ -1543,7 +1556,7 @@ tool_fini(void* /*tool_data*/) auto counters_output = counter_collection_buffered_output_t{tool::get_config().counter_collection}; auto scratch_memory_output = - scratch_memory_buffered_output_t{tool::get_config().scratch_memory}; + scratch_memory_buffered_output_t{tool::get_config().scratch_memory_trace}; auto node_id_sort = [](const auto& lhs, const auto& rhs) { return lhs.node_id < rhs.node_id; }; @@ -1561,7 +1574,7 @@ tool_fini(void* /*tool_data*/) rocprofiler::tool::generate_csv(tool_functions, _agents); } - auto contributions = std::unordered_map{}; + auto contributions = domain_stats_vec_t{}; generate_output(kernel_dispatch_output, contributions); generate_output(hsa_output, contributions); @@ -1580,6 +1593,7 @@ tool_fini(void* /*tool_data*/) { rocprofiler::tool::write_json(tool_functions, getpid(), + contributions, _agents, _counters, &hip_output.element_data, @@ -1617,6 +1631,11 @@ tool_fini(void* /*tool_data*/) &scratch_memory_output.element_data); } + if(tool::get_config().summary_output) + { + rocprofiler::tool::generate_stats(tool_functions, contributions); + } + auto destroy_output = [](auto& _buffered_output_v) { _buffered_output_v.destroy(); }; destroy_output(kernel_dispatch_output); diff --git a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp index 891de6bfc4..07f48a74b6 100644 --- a/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp +++ b/projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp @@ -72,7 +72,7 @@ struct async_copy_info; struct async_copy_info \ { \ static constexpr auto operation_idx = ROCPROFILER_MEMORY_COPY_##DIRECTION; \ - static constexpr auto name = #DIRECTION; \ + static constexpr auto name = "MEMORY_COPY_" #DIRECTION; \ }; SPECIALIZE_ASYNC_COPY_INFO(NONE) diff --git a/projects/rocprofiler-sdk/tests/bin/transpose/CMakeLists.txt b/projects/rocprofiler-sdk/tests/bin/transpose/CMakeLists.txt index f2bc686688..1781e6a2cc 100644 --- a/projects/rocprofiler-sdk/tests/bin/transpose/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/bin/transpose/CMakeLists.txt @@ -35,18 +35,65 @@ set(CMAKE_HIP_EXTENSIONS OFF) set(CMAKE_HIP_STANDARD_REQUIRED ON) set_source_files_properties(transpose.cpp PROPERTIES LANGUAGE HIP) -add_executable(transpose) -target_sources(transpose PRIVATE transpose.cpp) -target_compile_options(transpose PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -Werror) + +function(transpose_build_target _NAME _COMPILE_DEFS _LINK_TARGETS) + add_executable(${_NAME}) + target_sources(${_NAME} PRIVATE transpose.cpp) + target_compile_definitions(${_NAME} PRIVATE ${_COMPILE_DEFS}) + target_compile_options(${_NAME} PRIVATE -W -Wall -Wextra -Wpedantic -Wshadow -Werror) + target_link_libraries( + ${_NAME} PRIVATE Threads::Threads rocprofiler-sdk::tests-build-flags + ${_LINK_TARGETS}) + + if(TRANSPOSE_USE_MPI) + find_package(MPI REQUIRED) + target_compile_definitions(${_NAME} PRIVATE USE_MPI=1) + target_link_libraries(${_NAME} PRIVATE MPI::MPI_C) + endif() + +endfunction() find_package(Threads REQUIRED) -target_link_libraries(transpose PRIVATE Threads::Threads) - find_package(rocprofiler-sdk-roctx REQUIRED) -target_link_libraries(transpose PRIVATE rocprofiler-sdk-roctx::rocprofiler-sdk-roctx) -if(TRANSPOSE_USE_MPI) - find_package(MPI REQUIRED) - target_compile_definitions(transpose PRIVATE USE_MPI) - target_link_libraries(transpose PRIVATE MPI::MPI_C) +transpose_build_target(transpose "" rocprofiler-sdk-roctx::rocprofiler-sdk-roctx) + +# +# using old roctracer roctx +# +find_path( + roctracer_roctx_ROOT_DIR + NAMES include/roctracer/roctx.h + HINTS ${hip_DIR} + PATHS ${hip_DIR}) + +find_path( + roctracer_roctx_INCLUDE_DIR + NAMES roctracer/roctx.h + HINTS ${roctracer_roctx_ROOT_DIR} ${hip_DIR} + PATHS ${roctracer_roctx_ROOT_DIR} ${hip_DIR} + PATH_SUFFIXES include) + +find_library( + roctracer_roctx_LIBRARY + NAMES roctx64 + HINTS ${roctracer_roctx_ROOT_DIR} ${hip_DIR} + PATHS ${roctracer_roctx_ROOT_DIR} ${hip_DIR} + PATH_SUFFIXES lib lib64) + +include(FindPackageHandleStandardArgs) + +find_package_handle_standard_args(roctracer-roctx DEFAULT_MSG roctracer_roctx_ROOT_DIR + roctracer_roctx_INCLUDE_DIR roctracer_roctx_LIBRARY) + +if(roctracer-roctx_FOUND) + add_library(roctracer-roctx INTERFACE IMPORTED) + add_library(roctracer-roctx::roctracer-roctx ALIAS roctracer-roctx) + target_include_directories(roctracer-roctx INTERFACE ${roctracer_roctx_INCLUDE_DIR}) + target_link_libraries(roctracer-roctx INTERFACE ${roctracer_roctx_LIBRARY}) +endif() + +if(TARGET roctracer-roctx::roctracer-roctx) + transpose_build_target(transpose-roctracer-roctx "USE_ROCTRACER_ROCTX=1" + roctracer-roctx::roctracer-roctx) endif() diff --git a/projects/rocprofiler-sdk/tests/bin/transpose/transpose.cpp b/projects/rocprofiler-sdk/tests/bin/transpose/transpose.cpp index ad871d4c22..addb8bf938 100644 --- a/projects/rocprofiler-sdk/tests/bin/transpose/transpose.cpp +++ b/projects/rocprofiler-sdk/tests/bin/transpose/transpose.cpp @@ -20,8 +20,14 @@ // OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN // THE SOFTWARE. -#include "hip/hip_runtime.h" -#include "rocprofiler-sdk-roctx/roctx.h" +#include +#include + +#if defined(USE_ROCTRACER_ROCTX) +# include +#else +# include +#endif #include #include @@ -99,6 +105,15 @@ main(int argc, char** argv) printf("[transpose] Number of iterations: %zu\n", nitr); printf("[transpose] Syncing every %zu iterations\n", nsync); +#if defined(USE_ROCTRACER_ROCTX) + { + auto _roctracer_roctx_ss = std::stringstream{}; + _roctracer_roctx_ss << "roctracer/roctx v" << roctx_version_major() << "." + << roctx_version_minor(); + roctxMark(_roctracer_roctx_ss.str().c_str()); + } +#endif + #if defined(USE_MPI) MPI_Init(&argc, &argv); MPI_Comm_rank(MPI_COMM_WORLD, &rank); @@ -174,7 +189,15 @@ transpose(const int* in, int* out, int M, int N) void run(int rank, int tid, int devid, int argc, char** argv) { - roctxRangePush("run"); + auto roctx_run_id = roctxRangeStart("run"); + + const auto mark = [rank, tid, devid](std::string_view suffix) { + auto _ss = std::stringstream{}; + _ss << "run/rank-" << rank << "/thread-" << tid << "/device-" << devid << "/" << suffix; + roctxMark(_ss.str().c_str()); + }; + + mark("begin"); constexpr unsigned int M = 4960 * 2; constexpr unsigned int N = 4960 * 2; @@ -219,9 +242,16 @@ run(int rank, int tid, int devid, int argc, char** argv) auto t1 = std::chrono::high_resolution_clock::now(); for(size_t i = 0; i < nitr; ++i) { + roctxRangePush("run/iteration"); transpose<<>>(in, out, M, N); check_hip_error(); - if(i % nsync == (nsync - 1)) HIP_API_CALL(hipStreamSynchronize(stream)); + if(i % nsync == (nsync - 1)) + { + roctxRangePush("run/iteration/sync"); + HIP_API_CALL(hipStreamSynchronize(stream)); + roctxRangePop(); + } + roctxRangePop(); } auto t2 = std::chrono::high_resolution_clock::now(); HIP_API_CALL(hipStreamSynchronize(stream)); @@ -251,7 +281,9 @@ run(int rank, int tid, int devid, int argc, char** argv) delete[] inp_matrix; delete[] out_matrix; - roctxRangePop(); + mark("end"); + + roctxRangeStop(roctx_run_id); } namespace diff --git a/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py b/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py index d9751538d9..017c813055 100644 --- a/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py +++ b/projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py @@ -54,6 +54,16 @@ def test_otf2_data( otf2_data, json_data, categories=("hip", "hsa", "marker", "kernel", "memory_copy") ): + def get_operation_name(kind_id, op_id): + return json_data["rocprofiler-sdk-tool"]["strings"]["buffer_records"][kind_id][ + "operations" + ][op_id] + + def get_kind_name(kind_id): + return json_data["rocprofiler-sdk-tool"]["strings"]["buffer_records"][kind_id][ + "kind" + ] + mapping = { "hip": ("hip_api", "hip_api"), "hsa": ("hsa_api", "hsa_api"), @@ -72,6 +82,20 @@ def test_otf2_data( _otf2_data = otf2_data.loc[otf2_data["category"] == otf2_category] _json_data = json_data["rocprofiler-sdk-tool"]["buffer_records"][json_category] + # we do not encode the roctxMark "regions" in OTF2 because + # they don't map to the OTF2_REGION_ROLE_FUNCTION well + if json_category == "marker_api": + + def roctx_mark_filter(val): + return ( + None + if get_kind_name(val.kind) == "MARKER_CORE_API" + and get_operation_name(val.kind, val.operation) == "roctxMarkA" + else val + ) + + _json_data = [itr for itr in _json_data if roctx_mark_filter(itr) is not None] + assert len(_otf2_data) == len( _json_data ), f"{otf2_category} ({len(_otf2_data)}):\n\t{_otf2_data}\n{json_category} ({len(_json_data)}):\n\t{_json_data}" diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt index 334b84d1a4..24ba7aa51d 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt @@ -30,3 +30,4 @@ add_subdirectory(counter-collection) add_subdirectory(hsa-queue-dependency) add_subdirectory(kernel-rename) add_subdirectory(aborted-app) +add_subdirectory(summary) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-rename/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-rename/validate.py index c3b60f5ab3..609076b4f5 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-rename/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-rename/validate.py @@ -1,5 +1,6 @@ #!/usr/bin/env python3 +import re import sys import pytest @@ -83,7 +84,15 @@ def test_kernel_trace(json_data): def get_kind_name(kind_id): return data["strings"]["buffer_records"][kind_id]["kind"] - valid_kernel_names = ("run",) + valid_kernel_names = "|".join( + ( + "run", + "run/iteration", + "run/rank-([0-9]+)/thread-([0-9]+)/device-([0-9]+)/(begin|end)", + ) + ) + valid_kernel_regex = re.compile("^({})$".format(valid_kernel_names)) + kernel_dispatch_data = data["buffer_records"]["kernel_dispatch"] for dispatch in kernel_dispatch_data: assert get_kind_name(dispatch["kind"]) == "KERNEL_DISPATCH" @@ -97,6 +106,10 @@ def test_kernel_trace(json_data): assert dispatch["end_timestamp"] >= dispatch["start_timestamp"] kernel_name = get_kernel_name(dispatch_info["kernel_id"]) + assert ( + re.search(valid_kernel_regex, kernel_name) is None + ), f"kernel '{kernel_name}' matches regular expression '{valid_kernel_names}'" + assert kernel_name not in valid_kernel_names external_corr_id = dispatch["correlation_id"]["external"] @@ -104,7 +117,10 @@ def test_kernel_trace(json_data): kernel_rename = get_kernel_rename(external_corr_id) assert kernel_rename is not None, f"{dispatch}" - assert kernel_rename in valid_kernel_names + assert kernel_rename != kernel_name, f"{dispatch}" + assert ( + re.search(valid_kernel_regex, kernel_rename) is not None + ), f"renamed kernel '{kernel_rename}' does not match regular expression '{valid_kernel_names}'" def test_memory_copy_json_trace(json_data): diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/summary/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/summary/CMakeLists.txt new file mode 100644 index 0000000000..af386991a9 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/summary/CMakeLists.txt @@ -0,0 +1,127 @@ +# +# rocprofv3 tool tests for summaries +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-tests-rocprofv3-summary + LANGUAGES CXX + VERSION 0.0.0) + +find_package(rocprofiler-sdk REQUIRED) + +if(ROCPROFILER_MEMCHECK STREQUAL "LeakSanitizer") + set(LOG_LEVEL "warning") # info produces memory leak +else() + set(LOG_LEVEL "info") +endif() + +string(REPLACE "LD_PRELOAD=" "ROCPROF_PRELOAD=" PRELOAD_ENV + "${ROCPROFILER_MEMCHECK_PRELOAD_ENV}") + +set(summary-env "${PRELOAD_ENV}") + +rocprofiler_configure_pytest_files(CONFIG pytest.ini input-summary.yaml COPY validate.py + conftest.py) + +########################################################################################## +# +# Command line input +# +########################################################################################## + +set(summary-cmdl-env ${summary-env} "ARBITRARY_ENV_VARIABLE=%tag%-trace-cmd-line") + +add_test( + NAME rocprofv3-test-summary-cmd-line-execute + COMMAND + $ -d + ${CMAKE_CURRENT_BINARY_DIR}/%tag%-trace-cmd-line -o out --output-format pftrace + json otf2 --log-level env --runtime-trace --kernel-rename --summary + --summary-output-file summary --summary-per-domain --summary-groups + "KERNEL_DISPATCH|MEMORY_COPY" ".*_API" -- $ 2 500 10) + +set_tests_properties( + rocprofv3-test-summary-cmd-line-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${summary-cmdl-env}" + FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +add_test( + NAME rocprofv3-test-summary-cmd-line-validate + COMMAND + ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_results.json + --pftrace-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_results.pftrace + --otf2-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_results.otf2 + --summary-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_summary.txt) + +set(VALIDATION_FILES + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_results.pftrace + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_results.json + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-cmd-line/out_summary.txt) + +set_tests_properties( + rocprofv3-test-summary-cmd-line-validate + PROPERTIES TIMEOUT + 45 + LABELS + "integration-tests" + DEPENDS + "rocprofv3-test-summary-execute" + FAIL_REGULAR_EXPRESSION + "AssertionError" + ATTACHED_FILES_ON_FAIL + "${VALIDATION_FILES}") + +########################################################################################## +# +# YAML input +# +########################################################################################## + +set(summary-yaml-env ${summary-env} "ARBITRARY_ENV_VARIABLE=%tag%-trace-inp-yaml") + +add_test( + NAME rocprofv3-test-summary-inp-yaml-execute + COMMAND + $ -i + ${CMAKE_CURRENT_BINARY_DIR}/input-summary.yaml -- $ 2 500 + 10) + +set_tests_properties( + rocprofv3-test-summary-inp-yaml-execute + PROPERTIES TIMEOUT 45 LABELS "integration-tests" ENVIRONMENT "${summary-yaml-env}" + FAIL_REGULAR_EXPRESSION "${ROCPROFILER_DEFAULT_FAIL_REGEX}") + +add_test( + NAME rocprofv3-test-summary-inp-yaml-validate + COMMAND + ${Python3_EXECUTABLE} ${CMAKE_CURRENT_BINARY_DIR}/validate.py --json-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_results.json + --pftrace-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_results.pftrace + --otf2-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_results.otf2 + --summary-input + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_summary.txt) + +set(VALIDATION_FILES + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_results.pftrace + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_results.json + ${CMAKE_CURRENT_BINARY_DIR}/transpose-trace-inp-yaml/out_summary.txt) + +set_tests_properties( + rocprofv3-test-summary-inp-yaml-validate + PROPERTIES TIMEOUT + 45 + LABELS + "integration-tests" + DEPENDS + "rocprofv3-test-summary-inp-yaml-execute" + FAIL_REGULAR_EXPRESSION + "AssertionError" + ATTACHED_FILES_ON_FAIL + "${VALIDATION_FILES}") diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/summary/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/summary/conftest.py new file mode 100644 index 0000000000..f198d57371 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/summary/conftest.py @@ -0,0 +1,108 @@ +#!/usr/bin/env python3 + +import pandas as pd +import pytest +import json +import os +import io + +from rocprofiler_sdk.pytest_utils.dotdict import dotdict +from rocprofiler_sdk.pytest_utils import collapse_dict_list +from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader +from rocprofiler_sdk.pytest_utils.otf2_reader import OTF2Reader + + +def pytest_addoption(parser): + parser.addoption( + "--json-input", + action="store", + help="Path to JSON file.", + ) + parser.addoption( + "--pftrace-input", + action="store", + help="Path to Perfetto trace file.", + ) + parser.addoption( + "--otf2-input", + action="store", + help="Path to OTF2 trace file.", + ) + parser.addoption( + "--summary-input", + action="store", + help="Path to summary markdown file.", + ) + + +@pytest.fixture +def json_data(request): + filename = request.config.getoption("--json-input") + with open(filename, "r") as inp: + return dotdict(collapse_dict_list(json.load(inp))) + + +@pytest.fixture +def pftrace_data(request): + filename = request.config.getoption("--pftrace-input") + return PerfettoReader(filename).read()[0] + + +@pytest.fixture +def otf2_data(request): + filename = request.config.getoption("--otf2-input") + if not os.path.exists(filename): + raise FileExistsError(f"{filename} does not exist") + return OTF2Reader(filename).read()[0] + + +@pytest.fixture +def summary_data(request): + filename = request.config.getoption("--summary-input") + if not os.path.exists(filename): + raise FileExistsError(f"{filename} does not exist") + + domains = {} + with open(filename, "r") as inp: + lines = [itr.strip() for itr in inp.readlines()] + lines = [itr for itr in lines if itr and not itr.startswith("|--")] + + def rework(x): + tmp = [itr.strip(" ") for itr in x.split("|")] + tmp = [itr.strip() for itr in tmp if len(itr.strip()) > 0] + if tmp[0] == "NAME": + tmp = [f'"{itr}"' for itr in tmp] + else: + tmp[0] = f'"{tmp[0]}"' + tmp[1] = f'"{tmp[1]}"' + return ",".join(tmp) + + def process_current_domain(_name, _list): + if _name and _list: + _list = [rework(itr) for itr in _list] + _contents = "{}\n".format("\n".join(_list)) + ifs = io.StringIO(_contents) + df = pd.read_csv(ifs) + domains[_name] = df + _name = None + _list = [] + return (None, []) + + current_name = None + current_list = [] + + for itr in lines: + if not itr.startswith("|") or itr.startswith("ROCPROFV3 "): + current_name, current_list = process_current_domain( + current_name, current_list + ) + current_name = itr.strip().strip(":").replace("ROCPROFV3 ", "", 1) + rpos = current_name.rfind(" SUMMARY") + if rpos >= 0: + current_name = current_name[:rpos] + else: + current_list += [itr] + + process_current_domain(current_name, current_list) + + return domains diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/summary/input-summary.yaml b/projects/rocprofiler-sdk/tests/rocprofv3/summary/input-summary.yaml new file mode 100644 index 0000000000..c9568eda6c --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/summary/input-summary.yaml @@ -0,0 +1,11 @@ +jobs: + - output_directory: "@CMAKE_CURRENT_BINARY_DIR@/%env{ARBITRARY_ENV_VARIABLE}%" + output_file: out + output_format: [pftrace, json, otf2] + log_level: env + runtime_trace: True + kernel_rename: True + summary: True + summary_per_domain: True + summary_groups: ["KERNEL_DISPATCH|MEMORY_COPY"] + summary_output_file: "summary" diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/summary/pytest.ini b/projects/rocprofiler-sdk/tests/rocprofv3/summary/pytest.ini new file mode 100644 index 0000000000..5e1e1c14a0 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/summary/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s -vv +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/summary/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/summary/validate.py new file mode 100644 index 0000000000..526012630f --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocprofv3/summary/validate.py @@ -0,0 +1,267 @@ +#!/usr/bin/env python3 + +import os +import re +import sys +import pytest + + +def test_hip_api_trace(json_data): + data = json_data["rocprofiler-sdk-tool"] + + def get_operation_name(kind_id, op_id): + return data["strings"]["buffer_records"][kind_id]["operations"][op_id] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + valid_domain_names = ("HIP_RUNTIME_API",) + + hip_api_data = data["buffer_records"]["hip_api"] + + functions = [] + for api in hip_api_data: + kind = get_kind_name(api["kind"]) + assert kind in valid_domain_names + assert api["end_timestamp"] >= api["start_timestamp"] + functions.append(get_operation_name(api["kind"], api["operation"])) + + functions = list(set(functions)) + for itr in ( + "__hipPushCallConfiguration", + "__hipPopCallConfiguration", + "__hipRegisterFatBinary", + "__hipRegisterFunction", + ): + assert itr not in functions, f"{itr}" + + for itr in ( + "hipMallocAsync", + "hipMemcpyAsync", + "hipMemsetAsync", + "hipFreeAsync", + "hipLaunchKernel", + ): + assert itr in functions, f"{itr}" + + +def test_kernel_trace(json_data): + data = json_data["rocprofiler-sdk-tool"] + + def get_kernel_name(kernel_id): + return data["kernel_symbols"][kernel_id]["formatted_kernel_name"] + + def get_kernel_rename(corr_id): + for itr in data.strings.correlation_id.external: + if itr.key == corr_id: + return itr.value + return None + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + valid_kernel_names = "|".join( + ( + "run", + "run/iteration", + "run/rank-([0-9]+)/thread-([0-9]+)/device-([0-9]+)/(begin|end)", + ) + ) + valid_kernel_regex = re.compile("^({})$".format(valid_kernel_names)) + + kernel_dispatch_data = data["buffer_records"]["kernel_dispatch"] + for dispatch in kernel_dispatch_data: + assert get_kind_name(dispatch["kind"]) == "KERNEL_DISPATCH" + assert dispatch["correlation_id"]["internal"] > 0 + assert dispatch["correlation_id"]["external"] > 0 + + dispatch_info = dispatch["dispatch_info"] + assert dispatch_info["agent_id"]["handle"] > 0 + assert dispatch_info["queue_id"]["handle"] > 0 + assert dispatch_info["kernel_id"] > 0 + assert dispatch["end_timestamp"] >= dispatch["start_timestamp"] + + kernel_name = get_kernel_name(dispatch_info["kernel_id"]) + assert ( + re.search(valid_kernel_regex, kernel_name) is None + ), f"kernel '{kernel_name}' matches regular expression '{valid_kernel_names}'" + + assert kernel_name not in valid_kernel_names + + external_corr_id = dispatch["correlation_id"]["external"] + assert external_corr_id > 0 + + kernel_rename = get_kernel_rename(external_corr_id) + assert kernel_rename is not None, f"{dispatch}" + assert kernel_rename != kernel_name, f"{dispatch}" + assert ( + re.search(valid_kernel_regex, kernel_rename) is not None + ), f"renamed kernel '{kernel_rename}' does not match regular expression '{valid_kernel_names}'" + + +def test_memory_copy_trace(json_data): + data = json_data["rocprofiler-sdk-tool"] + + buffer_records = data["buffer_records"] + agent_data = data["agents"] + memory_copy_data = buffer_records["memory_copy"] + + def get_kind_name(kind_id): + return data["strings"]["buffer_records"][kind_id]["kind"] + + def get_agent(node_id): + for agent in agent_data: + if agent["id"]["handle"] == node_id["handle"]: + return agent + return None + + assert len(memory_copy_data) == 12 + + for row in memory_copy_data: + src_agent = get_agent(row["src_agent_id"]) + dst_agent = get_agent(row["dst_agent_id"]) + assert get_kind_name(row["kind"]) == "MEMORY_COPY" + assert src_agent is not None, f"{row}" + assert dst_agent is not None, f"{row}" + assert row["correlation_id"]["internal"] > 0 + assert row["end_timestamp"] >= row["start_timestamp"] + + +def test_metadata_data(json_data): + data = json_data["rocprofiler-sdk-tool"] + + # patch summary groups for testing purposes + # (it appears sometimes these contain single quotes and other times it doesn't) + data.metadata.config.summary_groups = [ + f"{itr}".strip("'") for itr in data.metadata.config.summary_groups + ] + + num_summary_grps = len(data.metadata.config.summary_groups) + expected_summary_grps = ( + [ + "KERNEL_DISPATCH|MEMORY_COPY", + ] + if num_summary_grps == 1 + else [ + "KERNEL_DISPATCH|MEMORY_COPY", + ".*_API", + ] + ) + + assert data.metadata.config.summary is True + assert data.metadata.config.summary_per_domain is True + assert data.metadata.config.summary_unit == "nsec" + assert data.metadata.config.summary_file == "summary" + assert data.metadata.config.summary_groups == expected_summary_grps + + assert len(data.metadata.command) == 4 + + # patch command to make it easier to test + data.metadata.command[0] = os.path.basename(data.metadata.command[0]) + + assert data.metadata.command == ["transpose", "2", "500", "10"] + + +def test_summary_data(json_data): + data = json_data["rocprofiler-sdk-tool"] + + domains = [] + for itr in data.summary: + domains.append(itr.domain) + if itr.domain == "KERNEL_DISPATCH": + assert itr.stats.count == 1004 + expected = dict([["run/iteration", 1000]]) + for oitr in itr.stats.operations: + if oitr.key in expected.keys(): + assert oitr.value.count == expected[oitr.key] + else: + assert oitr.key.startswith("run/rank-") + assert ( + re.match( + r"run/rank-([0-9]+)/thread-([0-9]+)/device-([0-9]+)/begin", + oitr.key, + ) + is not None + ) + assert oitr.value.count == 2 + elif itr.domain == "HIP_API": + assert itr.stats.count == 2135 + elif itr.domain == "MEMORY_COPY": + assert itr.stats.count == 12 + elif itr.domain == "MARKER_API": + assert itr.stats.count == 1106 + expected = dict( + [ + ["run", 2], + ["run/iteration", 1000], + ["run/iteration/sync", 100], + ] + ) + for oitr in itr.stats.operations: + if oitr.key in expected.keys(): + assert oitr.value.count == expected[oitr.key] + else: + assert oitr.key.startswith("run/rank-") + assert ( + re.match( + r"run/rank-([0-9]+)/thread-([0-9]+)/device-([0-9]+)/(begin|end)", + oitr.key, + ) + is not None + ) + assert oitr.value.count == 1 + else: + assert False, f"unhandled domain: {itr.domain}" + + assert len(list(set(domains))) == len(domains) + + +def test_summary_display_data(json_data, summary_data): + data = json_data["rocprofiler-sdk-tool"] + num_summary_grps = len(data.metadata.config.summary_groups) + + def get_df(domain): + return summary_data[domain] + + def get_dims(df): + # return rows x cols + return [df.shape[0], df.shape[1]] if df is not None else [0, 0] + + hip = get_df("HIP_API") + marker = get_df("MARKER_API") + dispatch = get_df("KERNEL_DISPATCH") + memcpy = get_df("MEMORY_COPY") + dispatch_and_copy = get_df("KERNEL_DISPATCH + MEMORY_COPY") + hip_and_marker = get_df("HIP_API + MARKER_API") if num_summary_grps > 1 else None + total = get_df("SUMMARY") + + expected_hip_and_marker_dims = [20, 9] if hip_and_marker is not None else [0, 0] + + assert get_dims(hip) == [13, 9] + assert get_dims(marker) == [7, 9] + assert get_dims(dispatch) == [3, 9] + assert get_dims(memcpy) == [2, 9] + assert get_dims(dispatch_and_copy) == [5, 9] + assert get_dims(hip_and_marker) == expected_hip_and_marker_dims + assert get_dims(total) == [22, 9] + + +def test_perfetto_data(pftrace_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + rocprofv3.test_perfetto_data( + pftrace_data, json_data, ("hip", "marker", "kernel", "memory_copy") + ) + + +def test_otf2_data(otf2_data, json_data): + import rocprofiler_sdk.tests.rocprofv3 as rocprofv3 + + rocprofv3.test_otf2_data( + otf2_data, json_data, ("hip", "marker", "kernel", "memory_copy") + ) + + +if __name__ == "__main__": + exit_code = pytest.main(["-x", __file__] + sys.argv[1:]) + sys.exit(exit_code) diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt index 660cd0e4dd..1de8ae9b88 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt @@ -53,8 +53,10 @@ add_test( ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_agent_info.csv --kernel-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv - --hip-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv - --hsa-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv + --hip-stats + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_api_stats.csv + --hsa-stats + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_api_stats.csv --memory-copy-stats ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv --json-input ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_results.json @@ -71,8 +73,8 @@ set(VALIDATION_FILES ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_trace.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_agent_info.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_kernel_stats.csv - ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_stats.csv - ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_stats.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hip_api_stats.csv + ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_hsa_api_stats.csv ${CMAKE_CURRENT_BINARY_DIR}/hip-in-libraries-trace/out_memory_copy_stats.csv) set_tests_properties( diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py index 600d77332d..d7fa202021 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py @@ -392,16 +392,19 @@ def test_memory_copy_trace( for row in memory_copy_input_data: assert row["Kind"] == "MEMORY_COPY" - assert row["Direction"] in ("HOST_TO_DEVICE", "DEVICE_TO_HOST") + assert row["Direction"] in ( + "MEMORY_COPY_HOST_TO_DEVICE", + "MEMORY_COPY_DEVICE_TO_HOST", + ) src_agent = get_agent(row["Source_Agent_Id"]) dst_agent = get_agent(row["Destination_Agent_Id"]) assert src_agent is not None and dst_agent is not None, f"{agent_info_input_data}" - if row["Direction"] == "HOST_TO_DEVICE": + if row["Direction"] == "MEMORY_COPY_HOST_TO_DEVICE": assert src_agent["Agent_Type"] == "CPU" assert dst_agent["Agent_Type"] == "GPU" - elif row["Direction"] == "DEVICE_TO_HOST": + elif row["Direction"] == "MEMORY_COPY_DEVICE_TO_HOST": assert src_agent["Agent_Type"] == "GPU" assert dst_agent["Agent_Type"] == "CPU" diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py index 014b9ae714..f107c39d34 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py @@ -170,13 +170,13 @@ def test_memory_copy_trace(agent_info_input_data, memory_copy_input_data): assert len(memory_copy_input_data) == 2 def test_row(idx, direction): - assert direction in ("HOST_TO_DEVICE", "DEVICE_TO_HOST") + assert direction in ("MEMORY_COPY_HOST_TO_DEVICE", "MEMORY_COPY_DEVICE_TO_HOST") row = memory_copy_input_data[idx] assert row["Direction"] == direction src_agent = get_agent(row["Source_Agent_Id"]) dst_agent = get_agent(row["Destination_Agent_Id"]) assert src_agent is not None and dst_agent is not None, f"{agent_info_input_data}" - if direction == "HOST_TO_DEVICE": + if direction == "MEMORY_COPY_HOST_TO_DEVICE": assert src_agent["Agent_Type"] == "CPU" assert dst_agent["Agent_Type"] == "GPU" else: @@ -185,8 +185,8 @@ def test_memory_copy_trace(agent_info_input_data, memory_copy_input_data): assert int(row["Correlation_Id"]) > 0 assert int(row["End_Timestamp"]) >= int(row["Start_Timestamp"]) - test_row(0, "HOST_TO_DEVICE") - test_row(1, "DEVICE_TO_HOST") + test_row(0, "MEMORY_COPY_HOST_TO_DEVICE") + test_row(1, "MEMORY_COPY_DEVICE_TO_HOST") def test_memory_copy_json_trace(json_data): @@ -208,14 +208,14 @@ def test_memory_copy_json_trace(json_data): assert len(memory_copy_data) == 2 def test_row(idx, direction): - assert direction in ("HOST_TO_DEVICE", "DEVICE_TO_HOST") + assert direction in ("MEMORY_COPY_HOST_TO_DEVICE", "MEMORY_COPY_DEVICE_TO_HOST") row = memory_copy_data[idx] src_agent = get_agent(row["src_agent_id"]) dst_agent = get_agent(row["dst_agent_id"]) assert get_kind_name(row["kind"]) == "MEMORY_COPY" assert src_agent is not None, f"{row}" assert dst_agent is not None, f"{row}" - if direction == "HOST_TO_DEVICE": + if direction == "MEMORY_COPY_HOST_TO_DEVICE": assert src_agent["type"] == 1 assert dst_agent["type"] == 2 else: @@ -224,8 +224,8 @@ def test_memory_copy_json_trace(json_data): assert row["correlation_id"]["internal"] > 0 assert row["end_timestamp"] >= row["start_timestamp"] - test_row(0, "HOST_TO_DEVICE") - test_row(1, "DEVICE_TO_HOST") + test_row(0, "MEMORY_COPY_HOST_TO_DEVICE") + test_row(1, "MEMORY_COPY_DEVICE_TO_HOST") def test_marker_api_trace(marker_input_data):