rocprofv3: summary reports + more JSON metadata (#1029)
* Move include/rocprofiler-sdk/cxx/details/delimit.hpp to tokenize.hpp
* Update docs/how-to/using-rocprofv3.rst
- fix code block indents
- reorder rocprofv3 options, limit them to important options
- add docs for `--runtime-trace`
* Update rocprofv3.py
- parser argument groups
- new `--runtime-trace` option
- new `--summary` option
- new `--summary-per-domain` option
- new `--summary-groups` option
- new `--summary-output-file` option
- new `--summary-units` option
* Update lib/rocprofiler-sdk/hsa/async_copy.cpp
- fix async copy operation names: add "MEMORY_COPY_" prefix
* lib/rocprofiler-sdk-tool: update statistics.{hpp,cpp}
- statistics<>::get_percent function
- stats_entry_t struct
- stats_formatter struct
- percentage struct
- std::to_string(::rocprofiler::tool::percentage)
* lib/rocprofiler-sdk-tool: update domain_type.{hpp,cpp}
- reorder domain_type enum values
* lib/rocprofiler-sdk-tool: update generateCSV.{hpp,cpp}
- separate writing CSV from accumulating statistics
- a lot of functionality was moved to statistics.{hpp,cpp}
* lib/rocprofiler-sdk-tool: update output_file.{hpp,cpp}
- output_stream_t struct
- get_output_stream(...) returns output_stream_t instance
* lib/rocprofiler-sdk-tool: update generateJSON.cpp
- update get_output_stream usage to output_stream_t
* lib/rocprofiler-sdk-tool: update generateOTF2.cpp
- header include order tweak
* lib/rocprofiler-sdk-tool: update buffered_output.hpp
- stats_data_t was renamed to stats_entry_t
* lib/rocprofiler-sdk-tool: update generatePerfetto.cpp
- header include tweak
* lib/rocprofiler-sdk-tool: update tmp_file_buffer.hpp
- emit warning message if write_ring_buffer fails after offloading instead of aborting
- prefer placement new instead of assignment in write_ring_buffer
* lib/rocprofiler-sdk-tool: add generateStats.{hpp,cpp}
- functions for accumulating statistics
* Update tests/rocprofv3/tracing-hip-in-libraries/CMakeLists.txt
- accommodate tweak to CSV output file name for HIP and HSA traces
* lib/rocprofiler-sdk-tool: update config.{hpp,cpp}
- new config variables
- stats_summary
- stats_summary_per_domain
- summary_output
- stats_summary_unit_value
- stats_summary_unit
- stats_summary_file
- stats_summary_groups
- support output keys for hostname: %hostname% / %h
* lib/rocprofiler-sdk-tool: update tool.cpp
- support summary output
* Documentation fixes
* Test for summary output
* Update tests/bin/transpose to use more ROCTx
- also support building with the roctracer ROCTx
* Remove roctxMark from OTF2 + fix kernel-rename tests
- following more ROCTx calls in transpose, kernel-rename validation had to be updated
* JSON metadata + JSON summary
- add serialization support for config
- add serialization support for statistics
- additions to json spec
- rocprofiler-sdk-tool/metadata/config
- rocprofiler-sdk-tool/metadata/command
- rocprofiler-sdk-tool/summary
- config output_keys support for NVIDIA %q{<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
This commit is contained in:
Jonathan R. Madsen
committed by
GitHub
GitHub
parent
1354452cf0
commit
395f01b689
@@ -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:
|
||||
|
||||
|
||||
Reference in New Issue
Block a user