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

rocprofv3: summary reports + more JSON metadata (#1029)

Browse Source 
* 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{<ENV-VAR>} 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: 395f01b689]
This commit is contained in:
Jonathan R. Madsen
2024-09-09 11:20:55 -05:00
committed by GitHub
parent 63f9bdc75e
commit 6743608258
40 changed files with 2293 additions and 580 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rocprofiler-sdk/source/bin/rocprofv3.py
+256 -111
View File
@@ -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,
projects/rocprofiler-sdk/source/docs/data/hip_stats.csv → projects/rocprofiler-sdk/source/docs/data/hip_api_stats.csv
View File
projects/rocprofiler-sdk/source/docs/data/hsa_trace.csv → projects/rocprofiler-sdk/source/docs/data/hsa_api_trace.csv
View File
projects/rocprofiler-sdk/source/docs/how-to/using-rocprofv3.rst
+183 -171
View File
@@ -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 <build-directory> <source-directory> -DCMAKE_PREFIX_PATH=/opt/rocm
cmake --build <build-directory> --target all --parallel <N>
cmake -B <build-directory> <source-directory> -DCMAKE_PREFIX_PATH=/opt/rocm
cmake --build <build-directory> --target all --parallel <N>
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 <tracing_option> -- <app_relative_path>
rocprofv3 <tracing_option> -- <app_relative_path>
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 -- <app_relative_path>
rocprofv3 -i input.txt -- <app_relative_path>
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 -- <app_relative_path>
rocprofv3 -i input.yml -- <app_relative_path>
$ 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:
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/CMakeLists.txt
+1 -1
View File
@@ -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}
projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/delimit.hpp → projects/rocprofiler-sdk/source/include/rocprofiler-sdk/cxx/details/tokenize.hpp
+58
View File
@@ -105,6 +105,64 @@ tokenize(
return _result;
}
/// \brief tokenize a string into a set
///
template <typename ContainerT = std::vector<std::string>,
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<DelimT>& 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<PredicateT>(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<PredicateT>(predicate)(line.substr(_start)));
break;
}
}
return _result;
}
/// \brief apply a string transformation to substring in between a common delimiter.
///
template <typename PredicateT = std::function<std::string(const std::string&)>>
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/CMakeLists.txt
+3
View File
@@ -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)
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/buffered_output.hpp
+1 -1
View File
@@ -59,7 +59,7 @@ struct buffered_output
operator bool() const { return enabled; }
std::deque<Tp> element_data = {};
stats_data_t stats = {};
stats_entry_t stats = {};
private:
bool enabled = false;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.cpp
+100 -7
View File
@@ -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 <rocprofiler-sdk/cxx/details/delimit.hpp>
#include <rocprofiler-sdk/cxx/details/tokenize.hpp>
#include <fmt/core.h>
#include <linux/limits.h>
#include <unistd.h>
#include <algorithm>
#include <chrono>
@@ -44,6 +46,7 @@
#include <set>
#include <sstream>
#include <string>
#include <string_view>
#include <vector>
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, 2>{std::regex{"(.*)%(env|ENV)\\{([A-Z0-9_]+)\\}%(.*)"},
std::regex{"(.*)\\$(env|ENV)\\{([A-Z0-9_]+)\\}(.*)"}};
new std::array<std::regex, 3>{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<char, PATH_MAX>{};
_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<std::string_view>{"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<std::string_view>{"##@@##"});
// 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<output_key>
@@ -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<output_key>{
{"%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<output_key>{
{"%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<output_key>& _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<char> 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<std::string>(_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<output_key>& _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)
{
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/config.hpp
+104 -16
View File
@@ -25,8 +25,14 @@
#include "lib/common/environment.hpp"
#include "lib/common/filesystem.hpp"
#include "lib/common/mpl.hpp"
#include <rocprofiler-sdk/cxx/serialization.hpp>
#include <fmt/format.h>
#include <set>
#include <sstream>
#include <string>
#include <unordered_set>
#include <vector>
@@ -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_context ContextT = config_context::global>
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<uint32_t> kernel_filter_range = {};
std::set<std::string> counters = {};
std::unordered_set<uint32_t> kernel_filter_range = {};
std::set<std::string> counters = {};
std::vector<std::string> stats_summary_groups = {};
template <typename ArchiveT>
void save(ArchiveT&) const;
template <typename ArchiveT>
void load(ArchiveT&)
{}
};
template <config_context ContextT = config_context::global>
template <typename ArchiveT>
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_context ContextT>
config&
get_config()
{
@@ -119,6 +204,11 @@ struct output_key
{
output_key(std::string _key, std::string _val, std::string _desc = {});
template <typename Tp,
typename Up = Tp,
std::enable_if_t<!common::mpl::is_string_type<Up>::value, int> = 0>
output_key(std::string _key, Tp&& _val, std::string _desc = {});
operator std::pair<std::string, std::string>() const;
std::string key = {};
@@ -126,16 +216,14 @@ struct output_key
std::string description = {};
};
template <typename Tp, typename Up, std::enable_if_t<!common::mpl::is_string_type<Up>::value, int>>
output_key::output_key(std::string _key, Tp&& _val, std::string _desc)
: key{std::move(_key)}
, value{fmt::format("{}", std::forward<Tp>(_val))}
, description{std::move(_desc)}
{}
std::vector<output_key>
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.cpp
+5 -5
View File
@@ -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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/domain_type.hpp
+3 -3
View File
@@ -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,
};
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.cpp
+77 -154
View File
@@ -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 <rocprofiler-sdk/fwd.h>
@@ -41,47 +42,6 @@ namespace tool
{
namespace
{
using stats_data_t = statistics<uint64_t, float_type>;
using stats_map_t = std::map<std::string_view, stats_data_t>;
struct percentage
{
float_type value = {};
friend std::ostream& operator<<(std::ostream& os, percentage val) { return (os << val.value); }
};
struct stats_formatter
{
template <typename Tp>
std::ostream& operator()(std::ostream& ofs, const Tp& _val) const
{
using value_type = common::mpl::unqualified_type_t<Tp>;
if constexpr(std::is_floating_point<value_type>::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<Tp, percentage>::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<std::pair<std::string_view, stats_data_t>>{};
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<rocprofiler_agent_v0_t>
}
}
stats_data_t
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& data)
const std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& 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<rocprofiler_buffer_tracing_hip_api_record_t>& data)
const std::deque<rocprofiler_buffer_tracing_hip_api_record_t>& 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<rocprofiler_buffer_tracing_hsa_api_record_t>& data)
const std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>& 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<rocprofiler_buffer_tracing_memory_copy_record_t>& data)
const std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>& 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<rocprofiler_buffer_tracing_marker_api_record_t>& data)
const std::deque<rocprofiler_buffer_tracing_marker_api_record_t>& 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<rocprofiler_tool_counter_collection_record_t>& data)
const std::deque<rocprofiler_tool_counter_collection_record_t>& 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<rocprofiler_buffer_tracing_scratch_memory_record_t>& data)
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& 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<domain_type, stats_data_t>& 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<stats_formatter>(_row,
@@ -688,9 +611,9 @@ generate_csv(tool_table* /*tool_functions*/, std::unordered_map<domain_type, sta
duration_ns,
avg_ns,
percentage{percent_v},
value.get_min(),
value.get_max(),
value.get_stddev());
value.total.get_min(),
value.total.get_max(),
value.total.get_stddev());
ofs << _row.str() << std::flush;
}
}
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateCSV.hpp
+35 -30
View File
@@ -22,6 +22,7 @@
#pragma once
#include "generateStats.hpp"
#include "helper.hpp"
#include "statistics.hpp"
@@ -31,41 +32,45 @@ namespace rocprofiler
{
namespace tool
{
using float_type = double;
using stats_data_t = statistics<uint64_t, float_type>;
void
generate_csv(tool_table* tool_functions, std::vector<rocprofiler_agent_v0_t>& data);
stats_data_t
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hip_api_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_marker_api_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_tool_counter_collection_record_t>& data);
stats_data_t
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& data);
const std::deque<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions, std::unordered_map<domain_type, stats_data_t>& data);
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hip_api_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_marker_api_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_tool_counter_collection_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& data,
const stats_entry_t& stats);
void
generate_csv(tool_table* tool_functions, const domain_stats_vec_t& data);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.cpp
+33 -8
View File
@@ -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 <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
@@ -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<rocprofiler_agent_v0_t> agent_data,
std::vector<rocprofiler_tool_counter_info_t> counter_data,
std::deque<rocprofiler_buffer_tracing_hip_api_record_t>* 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateJSON.hpp
+2
View File
@@ -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<rocprofiler_agent_v0_t> agent_data,
std::vector<rocprofiler_tool_counter_info_t> counter_data,
std::deque<rocprofiler_buffer_tracing_hip_api_record_t>* hip_api_deque,
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateOTF2.cpp
+7 -2
View File
@@ -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 <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
@@ -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<decltype(itr)>;
auto name = buffer_names.at(itr.kind, itr.operation);
auto paradigm = OTF2_PARADIGM_HIP;
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generatePerfetto.cpp
+7 -10
View File
@@ -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 <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
@@ -503,23 +503,20 @@ write_perfetto(
tracing_session->FlushBlocking();
tracing_session->StopBlocking();
auto filename = std::string{"results"};
auto cleanup = std::function<void(std::ostream*&)>{};
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<size_t>{0};
auto is_done = std::promise<void>{};
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<std::mutex>{_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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.cpp
+350
View File
@@ -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 <fmt/format.h>
#include <rocprofiler-sdk/fwd.h>
#include <rocprofiler-sdk/marker/api_id.h>
#include <unistd.h>
#include <cstdint>
#include <iomanip>
#include <regex>
#include <sstream>
#include <string_view>
#include <utility>
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<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& 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<rocprofiler_buffer_tracing_hip_api_record_t>& 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<rocprofiler_buffer_tracing_hsa_api_record_t>& 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<rocprofiler_buffer_tracing_memory_copy_record_t>& 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<rocprofiler_buffer_tracing_marker_api_record_t>& 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<rocprofiler_tool_counter_collection_record_t>& /*data*/)
{
return stats_entry_t{};
}
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& 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<std::string_view, domain_type>{};
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<double>(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<domain_stats_vec_t>{};
for(const auto& itr : tool::get_config().stats_summary_groups)
{
auto _names = std::vector<std::string>{};
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/generateStats.hpp
+63
View File
@@ -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<rocprofiler_buffer_tracing_kernel_dispatch_record_t>& data);
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hip_api_record_t>& data);
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_hsa_api_record_t>& data);
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_memory_copy_record_t>& data);
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_marker_api_record_t>& data);
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_tool_counter_collection_record_t>& data);
stats_entry_t
generate_stats(tool_table* tool_functions,
const std::deque<rocprofiler_buffer_tracing_scratch_memory_record_t>& data);
void
generate_stats(tool_table* tool_functions, const domain_stats_vec_t& data);
} // namespace tool
} // namespace rocprofiler
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.cpp
+13 -6
View File
@@ -33,6 +33,12 @@ namespace tool
{
namespace fs = common::filesystem;
namespace
{
const auto stdout_names = std::unordered_set<std::string_view>{"stdout", "STDOUT"};
const auto stderr_names = std::unordered_set<std::string_view>{"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<std::ostream*, output_stream_dtor_t>
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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/output_file.hpp
+45 -12
View File
@@ -26,7 +26,6 @@
#include "csv.hpp"
#include "lib/common/filesystem.hpp"
#include "lib/rocprofiler-sdk-tool/csv.hpp"
#include <array>
#include <iostream>
@@ -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<std::ostream*, ostream_dtor_t>;
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 <typename Tp>
std::ostream& operator<<(Tp&& value)
{
return ((stream) ? *stream : std::cerr) << std::forward<Tp>(value) << std::flush;
}
void close()
{
if(stream) (*stream) << std::flush;
if(dtor) dtor(stream);
}
bool writes_to_file() const { return (dynamic_cast<std::ofstream*>(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<std::ostream*, output_stream_dtor_t>
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<std::mutex>{m_mutex};
return ((m_stream) ? *m_stream : std::cerr) << std::forward<T>(value) << std::flush;
return ((m_os.stream) ? *m_os.stream : std::cerr) << std::forward<T>(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 <size_t N>
@@ -82,9 +116,8 @@ output_file::output_file(std::string name,
csv::csv_encoder<N> encoder,
std::array<std::string_view, N>&& 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.cpp
+75
View File
@@ -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 <typename Tp>
std::ostream&
stats_formatter::operator()(std::ostream& ofs, const Tp& _val) const
{
using value_type = common::mpl::unqualified_type_t<Tp>;
if constexpr(std::is_floating_point<value_type>::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<Tp, percentage>::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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/statistics.hpp
+120
View File
@@ -22,6 +22,13 @@
#pragma once
#include "domain_type.hpp"
#include "lib/common/logging.hpp"
#include "lib/common/mpl.hpp"
#include <rocprofiler-sdk/cxx/serialization.hpp>
#include <cmath>
#include <cstdint>
#include <fstream>
@@ -29,7 +36,9 @@
#include <iomanip>
#include <iostream>
#include <limits>
#include <map>
#include <type_traits>
#include <vector>
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 <typename ArchiveT>
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 Tp, typename Fp>
typename statistics<Tp, Fp>::float_type
statistics<Tp, Fp>::get_percent(float_type _total) const
{
constexpr float_type one_hundred = 100.0;
const float_type _sum = get_sum();
ROCP_WARNING_IF(static_cast<int64_t>(_sum) > static_cast<int64_t>(_total))
<< "percentage calculation > 100%. sum=" << _sum << " > total=" << _total;
return (_sum / _total) * one_hundred;
}
template <typename Tp, typename Fp>
typename statistics<Tp, Fp>::float_type
statistics<Tp, Fp>::get_percent(const statistics<Tp, Fp>& _rhs) const
{
return get_percent(_rhs.get_sum());
}
using float_type = double;
using stats_data_t = statistics<uint64_t, float_type>;
using stats_map_t = std::map<std::string_view, stats_data_t>;
using stats_pair_t = std::pair<std::string_view, stats_data_t>;
using stats_entry_vec_t = std::vector<stats_pair_t>;
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 <typename FuncT = sort_predicate_t>
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 <typename ArchiveT>
void serialize(ArchiveT& ar, const unsigned int) const
{
total.serialize(ar, 0);
auto _entries_map = std::map<std::string, stats_data_t>{};
for(const auto& itr : entries)
_entries_map.emplace(std::string{itr.first}, itr.second);
ar(cereal::make_nvp("operations", _entries_map));
}
};
template <typename FuncT>
stats_entry_t&
stats_entry_t::sort(FuncT&& _predicate)
{
std::sort(entries.begin(), entries.end(), std::forward<FuncT>(_predicate));
return *this;
}
using domain_stats_t = std::pair<domain_type, stats_entry_t>;
using domain_stats_vec_t = std::vector<domain_stats_t>;
struct stats_formatter
{
template <typename Tp>
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<Tp> 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
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tmp_file_buffer.hpp
+50 -3
View File
@@ -22,6 +22,7 @@
#pragma once
#include "domain_type.hpp"
#include "helper.hpp"
#include "tmp_file.hpp"
@@ -35,6 +36,7 @@
#include <mutex>
#include <string>
#include <tuple>
#include <type_traits>
#include <utility>
template <typename Tp>
@@ -71,15 +73,60 @@ void
write_ring_buffer(Tp _v, domain_type type)
{
auto [_tmp_buf, _tmp_file] = get_tmp_file_buffer<ring_buffer_t<Tp>>(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<ring_buffer_t<Tp>>(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<Tp>::value)
{
new(ptr) Tp{std::move(_v)};
}
else if constexpr(std::is_move_assignable<Tp>::value)
{
*ptr = std::move(_v);
}
else if constexpr(std::is_copy_constructible<Tp>::value)
{
new(ptr) Tp{_v};
}
else if constexpr(std::is_copy_assignable<Tp>::value)
{
*ptr = _v;
}
else
{
static_assert(std::is_void<Tp>::value,
"data type is neither move/copy constructible nor move/copy assignable");
}
}
*ptr = std::move(_v);
}
template <typename Tp>
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk-tool/tool.cpp
+26 -7
View File
@@ -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 <typename Tp, domain_type DomainT>
void
generate_output(rocprofiler::tool::buffered_output<Tp, DomainT>& output_v,
std::unordered_map<domain_type, stats_data_t>& 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<domain_type, stats_data_t>{};
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);
projects/rocprofiler-sdk/source/lib/rocprofiler-sdk/hsa/async_copy.cpp
+1 -1
View File
@@ -72,7 +72,7 @@ struct async_copy_info;
struct async_copy_info<ROCPROFILER_MEMORY_COPY_##DIRECTION> \
{ \
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)
projects/rocprofiler-sdk/tests/bin/transpose/CMakeLists.txt
+57 -10
View File
@@ -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()
projects/rocprofiler-sdk/tests/bin/transpose/transpose.cpp
+37 -5
View File
@@ -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 <hip/hip_runtime.h>
#include <sstream>
#if defined(USE_ROCTRACER_ROCTX)
# include <roctracer/roctx.h>
#else
# include <rocprofiler-sdk-roctx/roctx.h>
#endif
#include <chrono>
#include <cstdio>
@@ -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<<<grid, block, 0, stream>>>(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
projects/rocprofiler-sdk/tests/pytest-packages/tests/rocprofv3.py
+24
View File
@@ -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}"
projects/rocprofiler-sdk/tests/rocprofv3/CMakeLists.txt
+1
View File
@@ -30,3 +30,4 @@ add_subdirectory(counter-collection)
add_subdirectory(hsa-queue-dependency)
add_subdirectory(kernel-rename)
add_subdirectory(aborted-app)
add_subdirectory(summary)
projects/rocprofiler-sdk/tests/rocprofv3/kernel-rename/validate.py
+18 -2
View File
@@ -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):
projects/rocprofiler-sdk/tests/rocprofv3/summary/CMakeLists.txt
+127
View File
@@ -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
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> -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" -- $<TARGET_FILE:transpose> 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
$<TARGET_FILE:rocprofiler-sdk::rocprofv3> -i
${CMAKE_CURRENT_BINARY_DIR}/input-summary.yaml -- $<TARGET_FILE:transpose> 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}")
projects/rocprofiler-sdk/tests/rocprofv3/summary/conftest.py
+108
View File
@@ -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
projects/rocprofiler-sdk/tests/rocprofv3/summary/input-summary.yaml
+11
View File
@@ -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"
projects/rocprofiler-sdk/tests/rocprofv3/summary/pytest.ini
+5
View File
@@ -0,0 +1,5 @@
[pytest]
addopts = --durations=20 -rA -s -vv
testpaths = validate.py
pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages
projects/rocprofiler-sdk/tests/rocprofv3/summary/validate.py
+267
View File
@@ -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)
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt
+6 -4
View File
@@ -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(
projects/rocprofiler-sdk/tests/rocprofv3/tracing-hip-in-libraries/validate.py
+6 -3
View File
@@ -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"
projects/rocprofiler-sdk/tests/rocprofv3/tracing/validate.py
+8 -8
View File
@@ -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):