Add 'projects/rocprofiler-systems/' from commit '92e1d84c72c9321d79a1866e0090fae0215e6557'

git-subtree-dir: projects/rocprofiler-systems
git-subtree-mainline: ee9e74df21
git-subtree-split: 92e1d84c72

projects/rocprofiler-systems/docs/.gitignore

@@ -0,0 +1,2 @@
+_build/
+_doxygen/

projects/rocprofiler-systems/docs/conceptual/data-collection-modes.rst

@@ -0,0 +1,147 @@
+.. meta::
+   :description: ROCm Systems Profiler data collection modes documentation
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, data collection, tracking, visualization, tool, Instinct, accelerator, AMD
+
+**********************
+Data collection modes
+**********************
+
+ROCm Systems Profiler supports several modes of recording trace and profiling data for your application.
+
+.. note::
+
+   For an explanation of the terms used in this topic, see
+   the :doc:`ROCm Systems Profiler glossary <../reference/rocprof-sys-glossary>`.
+
++-----------------------------+---------------------------------------------------------+
+| Mode                        | Description                                             |
++=============================+=========================================================+
+| Binary Instrumentation      | Locates functions (and loops, if desired) in the binary |
+|                             | and inserts snippets at the entry and exit              |
++-----------------------------+---------------------------------------------------------+
+| Statistical Sampling        | Periodically pauses application at specified intervals  |
+|                             | and records various metrics for the given call stack    |
++-----------------------------+---------------------------------------------------------+
+| Callback APIs               | Parallelism frameworks such as ROCm, OpenMP, and Kokkos |
+|                             | make callbacks into ROCm Systems Profiler to provide    |
+|                             | information about the work the API is performing        |
++-----------------------------+---------------------------------------------------------+
+| Dynamic Symbol Interception | Wrap function symbols defined in a position independent |
+|                             | dynamic library/executable, like ``pthread_mutex_lock`` |
+|                             | in ``libpthread.so`` or ``MPI_Init`` in the MPI library |
++-----------------------------+---------------------------------------------------------+
+| User API                    | User-defined regions and controls for ROCm Systems      |
+|                             | Profiler                                                |
++-----------------------------+---------------------------------------------------------+
+
+The two most generic and important modes are binary instrumentation and statistical sampling.
+It is important to understand their advantages and disadvantages.
+Binary instrumentation and statistical sampling can be performed with the ``rocprof-sys-instrument``
+executable. For statistical sampling, it's highly recommended to use the
+``rocprof-sys-sample`` executable instead if binary instrumentation isn't required or needed.
+Callback APIs and dynamic symbol interception can be utilized with either tool.
+
+Binary instrumentation
+-----------------------------------
+
+Binary instrumentation lets you record deterministic measurements for
+every single invocation of a given function.
+Binary instrumentation effectively adds instructions to the target application to
+collect the required information. It therefore has the potential to cause performance
+changes which might, in some cases, lead to inaccurate results. The effect depends on
+the information being collected and which features are activated in ROCm Systems Profiler.
+For example, collecting only the wall-clock timing data
+has less of an effect than collecting the wall-clock timing, CPU-clock timing,
+memory usage, cache-misses, and number of instructions that were run. Similarly,
+collecting a flat profile has less overhead than a hierarchical profile
+and collecting a trace OR a profile has less overhead than collecting a
+trace AND a profile.
+
+In ROCm Systems Profiler, the primary heuristic for controlling the overhead with binary
+instrumentation is the minimum number of instructions for selecting functions
+for instrumentation.
+
+Statistical sampling
+-----------------------------------
+
+Statistical call-stack sampling periodically interrupts the application at
+regular intervals using operating system interrupts.
+Sampling is typically less numerically accurate and specific, but the
+target program runs at nearly full speed.
+In contrast to the data derived from binary instrumentation, the resulting
+data is not exact but is instead a statistical approximation.
+However, sampling often provides a more accurate picture of the application
+execution because it is less intrusive to the target application and has fewer
+side effects on memory caches or instruction decoding pipelines. Furthermore,
+because sampling does not affect the execution speed as much, is it
+relatively immune to over-evaluating the cost of small, frequently called
+functions or "tight" loops.
+
+In ROCm Systems Profiler, the overhead for statistical sampling depends on the
+sampling rate and whether the samples are taken with respect to the CPU time
+and/or real time.
+
+Binary instrumentation vs. statistical sampling example
+-------------------------------------------------------
+
+Consider the following code:
+
+.. code-block:: c++
+
+   long fib(long n)
+   {
+        if(n < 2) return n;
+        return fib(n - 1) + fib(n - 2);
+   }
+
+   void run(long n)
+   {
+        long result = fib(n);
+        printf("[%li] fibonacci(%li) = %li\n", i, n, result);
+   }
+
+   int main(int argc, char** argv)
+   {
+        long nfib = 30;
+        long nitr = 10;
+        if(argc > 1) nfib = atol(argv[1]);
+        if(argc > 2) nitr = atol(argv[2]);
+
+        for(long i = 0; i < nitr; ++i)
+            run(nfib);
+
+        return 0;
+   }
+
+Binary instrumentation of the ``fib`` function will record **every single invocation**
+of the function. For a very small function
+such as ``fib``, this results in **significant** overhead since this simple function
+takes about 20 instructions, whereas the entry and
+exit snippets are ~1024 instructions. Therefore, you generally want to avoid
+instrumenting functions where the instrumented function has significantly fewer
+instructions than entry and exit instrumentation. (Note that many of the
+instructions in entry and exit functions are either logging functions or
+depend on the runtime settings and thus might never run). However,
+due to the number of potential instructions in the entry and exit snippets,
+the default behavior of ``rocprof-sys-instrument`` is to only instrument functions
+which contain at least 1024 instructions.
+
+However, recording every single invocation of the function can be extremely
+useful for detecting anomalies, such as profiles that show minimum or maximum values much smaller or larger
+than the average or a high standard deviation. In this case, the traces help you
+identify exactly when and where those instances deviated from the norm.
+Compare the level of detail in the following traces. In the top image,
+every instance of the ``fib`` function is instrumented, while in the bottom image,
+the ``fib`` call-stack is derived via sampling.
+
+Binary instrumentation of the Fibonacci function
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+.. image:: ../data/fibonacci-instrumented.png
+   :alt: Visualization of the output of a binary instrumentation of the Fibonacci function
+
+Statistical sampling of the Fibonacci function
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+.. image:: ../data/fibonacci-sampling.png
+   :alt: Visualization of the output of a statistical sample of the Fibonacci function

projects/rocprofiler-systems/docs/conceptual/rocprof-sys-feature-set.rst

@@ -0,0 +1,137 @@
+.. meta::
+   :description: ROCm Systems Profiler feature set documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, feature set, use cases, tracking, visualization, tool, Instinct, accelerator, AMD
+
+********************************************
+ROCm Systems Profiler features and use cases
+********************************************
+
+`ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ is designed to be highly extensible.
+Internally, it leverages the `Timemory performance analysis toolkit <https://github.com/ROCm/timemory>`_
+to manage extensions, resources, data, and other items. It supports the following features,
+modes, metrics, and APIs.
+
+Data collection modes
+========================================
+
+* Dynamic instrumentation
+
+  * Runtime instrumentation: Instrument executables and shared libraries at runtime
+  * Binary rewriting: Generate a new executable and/or library with instrumentation built-in
+
+* Statistical sampling: Periodic software interrupts per-thread
+* Process-level sampling: A background thread records process-, system- and device-level metrics while the application runs
+* Causal profiling: Quantifies the potential impact of optimizations in parallel code
+
+Data analysis
+========================================
+
+* High-level summary profiles with mean, min, max, and standard deviation statistics
+
+  * Low overhead and memory efficient
+  * Ideal for running at scale
+
+* Comprehensive traces for every individual event and measurement
+* Application speed-up predictions resulting from potential optimizations in functions and lines of code based on causal profiling
+
+Parallelism API support
+========================================
+
+* HIP
+* HSA
+* Pthreads
+* MPI
+* Kokkos-Tools (KokkosP)
+* OpenMP-Tools (OMPT)
+
+GPU metrics
+========================================
+
+* GPU hardware counters
+* HIP API tracing
+* HIP kernel tracing
+* HSA API tracing
+* HSA operation tracing
+* rocDecode API tracing
+* rocJPEG API tracing
+* System-level sampling (via AMD-SMI)
+
+  * Memory usage
+  * Power usage
+  * Temperature
+  * Utilization
+  * VCN activity
+  * JPEG activity
+    Note: The availability of VCN and JPEG engine activity depends on device support for different ASICs. If unsupported, all values for VCN_ACTIVITY and JPEG_ACTIVITY will be reported as N/A in the output of amd-smi metric--usage.
+
+CPU metrics
+========================================
+
+* CPU hardware counters sampling and profiles
+* CPU frequency sampling
+* Various timing metrics
+
+  * Wall time
+  * CPU time (process and thread)
+  * CPU utilization (process and thread)
+  * User CPU time
+  * Kernel CPU time
+
+* Various memory metrics
+
+  * High-water mark (sampling and profiles)
+  * Memory page allocation
+  * Virtual memory usage
+
+* Network statistics
+* I/O metrics
+* Many others
+
+Third-party API support
+========================================
+
+* TAU
+* LIKWID
+* Caliper
+* CrayPAT
+* VTune
+* NVTX
+* ROCTX
+
+ROCm Systems Profiler use cases
+========================================
+
+When analyzing the performance of an application, do NOT
+assume you know where the performance bottlenecks are
+and why they are happening. ROCm Systems Profiler is a tool for analyzing the entire
+application and its performance. It is
+ideal for characterizing where optimization would have the greatest impact
+on an end-to-end run of the application and for
+viewing what else is happening on the system during a performance bottleneck.
+
+When GPUs are involved, there is a tendency to assume that
+the quickest path to performance improvement is minimizing
+the runtime of the GPU kernels. This is a highly flawed assumption.
+If you optimize the runtime of a kernel from one millisecond
+to 1 microsecond (1000x speed-up) but the original application never
+spent time waiting for kernels to complete,
+there would be no statistically significant reduction in the end-to-end
+runtime of your application. In other words, it does not matter
+how fast or slow the code on GPU is if the application has a
+bottleneck on waiting on the GPU.
+
+Use ROCm Systems Profiler to obtain a high-level view of the entire application. Use it
+to determine where the performance bottlenecks are and
+obtain clues to why these bottlenecks are happening. Rather than worrying about kernel
+performance, start your investigation with ROCm Systems Profiler, which characterizes the
+broad picture.
+
+.. note::
+
+   For insight into the execution of individual kernels on the GPU,
+   use `ROCm Compute Profiler <https://github.com/rocm/rocprofiler-compute>`_.
+
+In terms of CPU analysis, ROCm Systems Profiler does not target any specific vendor.
+It works just as well on AMD and non-AMD CPUs.
+With regard to the GPU, ROCm Systems Profiler is currently restricted to HIP and HSA APIs
+and kernels running on AMD GPUs.

projects/rocprofiler-systems/docs/conf.py

@@ -0,0 +1,56 @@
+# MIT License
+
+# Copyright (c) 2023 - 2025 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.
+
+# Configuration file for the Sphinx documentation builder.
+#
+# This file only contains a selection of the most common options. For a full
+# list see the documentation:
+# https://www.sphinx-doc.org/en/master/usage/configuration.html
+
+import re
+
+from rocm_docs import ROCmDocs
+
+with open("../VERSION", encoding="utf-8") as f:
+    match = re.search(r"([0-9.]+)[^0-9.]+", f.read())
+    if not match:
+        raise ValueError("VERSION not found!")
+    version_number = match[1]
+
+external_projects_current_project = "rocprofiler-systems"
+
+project = "rocprofiler-systems"
+author = "Advanced Micro Devices, Inc."
+copyright = "Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved."
+version = version_number
+release = version_number
+html_title = f"ROCm Systems Profiler {version} documentation"
+
+external_toc_path = "./sphinx/_toc.yml"
+
+docs_core = ROCmDocs(html_title)
+docs_core.setup()
+docs_core.run_doxygen(doxygen_root="doxygen", doxygen_path="doxygen/xml")
+docs_core.enable_api_reference()
+
+for sphinx_var in ROCmDocs.SPHINX_VARS:
+    globals()[sphinx_var] = getattr(docs_core, sphinx_var)

projects/rocprofiler-systems/docs/doxygen/.gitignore

@@ -0,0 +1,3 @@
+html/
+latex/
+xml/

projects/rocprofiler-systems/docs/doxygen/Doxyfile

@@ -0,0 +1,373 @@
+# Doxyfile 1.8.20
+
+#---------------------------------------------------------------------------
+# Project related configuration options
+#---------------------------------------------------------------------------
+DOXYFILE_ENCODING      = UTF-8
+PROJECT_NAME           = rocprofiler-systems
+PROJECT_NUMBER         = 1.11.3
+PROJECT_BRIEF          = "High-level and comprehensive application tracing and profiling on both the CPU and GPU"
+PROJECT_LOGO           =
+OUTPUT_DIRECTORY       = .
+CREATE_SUBDIRS         = NO
+ALLOW_UNICODE_NAMES    = YES
+OUTPUT_LANGUAGE        = English
+OUTPUT_TEXT_DIRECTION  = None
+BRIEF_MEMBER_DESC      = YES
+REPEAT_BRIEF           = YES
+ABBREVIATE_BRIEF       =
+ALWAYS_DETAILED_SEC    = YES
+INLINE_INHERITED_MEMB  = YES
+FULL_PATH_NAMES        = YES
+STRIP_FROM_PATH        = /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprofiler-systems/checkouts/
+STRIP_FROM_INC_PATH    = /home/docs/checkouts/readthedocs.org/user_builds/advanced-micro-devices-rocprofiler-systems/checkouts/
+SHORT_NAMES            = NO
+JAVADOC_AUTOBRIEF      = NO
+JAVADOC_BANNER         = NO
+QT_AUTOBRIEF           = NO
+MULTILINE_CPP_IS_BRIEF = YES
+PYTHON_DOCSTRING       = YES
+INHERIT_DOCS           = YES
+SEPARATE_MEMBER_PAGES  = NO
+TAB_SIZE               = 4
+ALIASES                =
+OPTIMIZE_OUTPUT_FOR_C  = NO
+OPTIMIZE_OUTPUT_JAVA   = NO
+OPTIMIZE_FOR_FORTRAN   = NO
+OPTIMIZE_OUTPUT_VHDL   = NO
+OPTIMIZE_OUTPUT_SLICE  = NO
+EXTENSION_MAPPING      = hpp=C++ \
+                         cpp=C++ \
+                         hh=C++ \
+                         cc=C++ \
+                         h=C \
+                         c=C \
+                         py=Python
+MARKDOWN_SUPPORT       = YES
+TOC_INCLUDE_HEADINGS   = 2
+AUTOLINK_SUPPORT       = YES
+BUILTIN_STL_SUPPORT    = YES
+CPP_CLI_SUPPORT        = NO
+SIP_SUPPORT            = NO
+IDL_PROPERTY_SUPPORT   = YES
+DISTRIBUTE_GROUP_DOC   = NO
+GROUP_NESTED_COMPOUNDS = YES
+SUBGROUPING            = YES
+INLINE_GROUPED_CLASSES = NO
+INLINE_SIMPLE_STRUCTS  = YES
+TYPEDEF_HIDES_STRUCT   = NO
+LOOKUP_CACHE_SIZE      = 5
+NUM_PROC_THREADS       = 0
+#---------------------------------------------------------------------------
+# Build related configuration options
+#---------------------------------------------------------------------------
+EXTRACT_ALL            = YES
+EXTRACT_PRIVATE        = NO
+EXTRACT_PRIV_VIRTUAL   = NO
+EXTRACT_PACKAGE        = NO
+EXTRACT_STATIC         = NO
+EXTRACT_LOCAL_CLASSES  = YES
+EXTRACT_LOCAL_METHODS  = NO
+EXTRACT_ANON_NSPACES   = NO
+HIDE_UNDOC_MEMBERS     = NO
+HIDE_UNDOC_CLASSES     = YES
+HIDE_FRIEND_COMPOUNDS  = NO
+HIDE_IN_BODY_DOCS      = NO
+INTERNAL_DOCS          = NO
+CASE_SENSE_NAMES       = NO
+HIDE_SCOPE_NAMES       = NO
+HIDE_COMPOUND_REFERENCE= NO
+SHOW_INCLUDE_FILES     = YES
+SHOW_GROUPED_MEMB_INC  = NO
+FORCE_LOCAL_INCLUDES   = YES
+INLINE_INFO            = YES
+SORT_MEMBER_DOCS       = YES
+SORT_BRIEF_DOCS        = NO
+SORT_MEMBERS_CTORS_1ST = YES
+SORT_GROUP_NAMES       = NO
+SORT_BY_SCOPE_NAME     = NO
+STRICT_PROTO_MATCHING  = NO
+GENERATE_TODOLIST      = NO
+GENERATE_TESTLIST      = NO
+GENERATE_BUGLIST       = NO
+GENERATE_DEPRECATEDLIST= NO
+ENABLED_SECTIONS       =
+MAX_INITIALIZER_LINES  = 30
+SHOW_USED_FILES        = YES
+SHOW_FILES             = YES
+SHOW_NAMESPACES        = YES
+FILE_VERSION_FILTER    =
+LAYOUT_FILE            =
+CITE_BIB_FILES         =
+#---------------------------------------------------------------------------
+# Configuration options related to warning and progress messages
+#---------------------------------------------------------------------------
+QUIET                  = NO
+WARNINGS               = YES
+WARN_IF_UNDOCUMENTED   = YES
+WARN_IF_DOC_ERROR      = YES
+WARN_NO_PARAMDOC       = YES
+WARN_AS_ERROR          = YES
+WARN_FORMAT            = "---> WARNING!   $file:$line: $text"
+WARN_LOGFILE           = doc/warnings.log
+#---------------------------------------------------------------------------
+# Configuration options related to the input files
+#---------------------------------------------------------------------------
+INPUT                  = ../../README.md \
+                         ../../source/lib/rocprof-sys-user/rocprofiler-systems/types.h \
+                         ../../source/lib/rocprof-sys-user/rocprofiler-systems/categories.h \
+                         ../../source/lib/rocprof-sys-user/rocprofiler-systems/user.h \
+                         ../../source/lib/rocprof-sys-user/rocprofiler-systems/causal.h
+INPUT_ENCODING         = UTF-8
+FILE_PATTERNS          = *.h \
+                         *.hh \
+                         *.hpp \
+                         *.c \
+                         *.cc \
+                         *.cxx \
+                         *.cpp \
+                         *.c++ \
+                         *.icc \
+                         *.tcc \
+                         *.py
+RECURSIVE              = YES
+EXCLUDE                =
+EXCLUDE_SYMLINKS       = YES
+EXCLUDE_PATTERNS       = */.git/* \
+                         ../../external/* \
+                         ../../examples/* \
+                         ../../tests/*
+EXCLUDE_SYMBOLS        = "std::*" \
+                         "ROCPROFSYS_ATTRIBUTE" \
+                         "ROCPROFSYS_VISIBILITY" \
+                         "ROCPROFSYS_PUBLIC_API" \
+                         "ROCPROFSYS_HIDDEN_API" \
+                         "SpaceHandle" \
+                         "KokkosPDevice*"
+EXAMPLE_PATH           = ../../examples
+EXAMPLE_PATTERNS       = *.h \
+                         *.hh \
+                         *.hpp \
+                         *.c \
+                         *.cc \
+                         *.cpp \
+                         *.py \
+                         *.txt
+EXAMPLE_RECURSIVE      = YES
+IMAGE_PATH             =
+INPUT_FILTER           =
+FILTER_PATTERNS        =
+FILTER_SOURCE_FILES    = NO
+FILTER_SOURCE_PATTERNS =
+USE_MDFILE_AS_MAINPAGE = ../../README.md
+#---------------------------------------------------------------------------
+# Configuration options related to source browsing
+#---------------------------------------------------------------------------
+SOURCE_BROWSER         = YES
+INLINE_SOURCES         = YES
+STRIP_CODE_COMMENTS    = NO
+REFERENCED_BY_RELATION = YES
+REFERENCES_RELATION    = YES
+REFERENCES_LINK_SOURCE = YES
+SOURCE_TOOLTIPS        = YES
+USE_HTAGS              = NO
+VERBATIM_HEADERS       = YES
+#---------------------------------------------------------------------------
+# Configuration options related to the alphabetical class index
+#---------------------------------------------------------------------------
+ALPHABETICAL_INDEX     = YES
+COLS_IN_ALPHA_INDEX    = 5
+IGNORE_PREFIX          =
+#---------------------------------------------------------------------------
+# Configuration options related to the HTML output
+#---------------------------------------------------------------------------
+GENERATE_HTML          = YES
+HTML_OUTPUT            = html
+HTML_FILE_EXTENSION    = .html
+HTML_HEADER            = ../_doxygen/header.html
+HTML_FOOTER            = ../_doxygen/footer.html
+HTML_STYLESHEET        = ../_doxygen/stylesheet.css
+HTML_EXTRA_STYLESHEET  = ../_doxygen/extra_stylesheet.css
+HTML_EXTRA_FILES       =
+HTML_COLORSTYLE_HUE    = 220
+HTML_COLORSTYLE_SAT    = 100
+HTML_COLORSTYLE_GAMMA  = 80
+HTML_TIMESTAMP         = YES
+HTML_DYNAMIC_MENUS     = YES
+HTML_DYNAMIC_SECTIONS  = YES
+HTML_INDEX_NUM_ENTRIES = 1000
+GENERATE_DOCSET        = NO
+DOCSET_FEEDNAME        = "Doxygen generated docs"
+DOCSET_BUNDLE_ID       = org.doxygen.rocprofiler-systems
+DOCSET_PUBLISHER_ID    = org.doxygen.amd
+DOCSET_PUBLISHER_NAME  = "Advanced Micro Devices, Inc."
+GENERATE_HTMLHELP      = NO
+CHM_FILE               =
+HHC_LOCATION           =
+GENERATE_CHI           = NO
+CHM_INDEX_ENCODING     =
+BINARY_TOC             = NO
+TOC_EXPAND             = YES
+GENERATE_QHP           = NO
+QCH_FILE               =
+QHP_NAMESPACE          =
+QHP_VIRTUAL_FOLDER     = doc
+QHP_CUST_FILTER_NAME   =
+QHP_CUST_FILTER_ATTRS  =
+QHP_SECT_FILTER_ATTRS  =
+QHG_LOCATION           =
+GENERATE_ECLIPSEHELP   = NO
+ECLIPSE_DOC_ID         = org.doxygen.rocprofiler-systems
+DISABLE_INDEX          = NO
+GENERATE_TREEVIEW      = NO
+ENUM_VALUES_PER_LINE   = 1
+TREEVIEW_WIDTH         = 300
+EXT_LINKS_IN_WINDOW    = YES
+HTML_FORMULA_FORMAT    = png
+FORMULA_FONTSIZE       = 12
+FORMULA_TRANSPARENT    = YES
+FORMULA_MACROFILE      =
+USE_MATHJAX            = NO
+MATHJAX_FORMAT         = HTML-CSS
+MATHJAX_RELPATH        = http://cdn.mathjax.org/mathjax/latest
+MATHJAX_EXTENSIONS     =
+MATHJAX_CODEFILE       =
+SEARCHENGINE           = NO
+SERVER_BASED_SEARCH    = NO
+EXTERNAL_SEARCH        = NO
+SEARCHENGINE_URL       =
+SEARCHDATA_FILE        = searchdata.xml
+EXTERNAL_SEARCH_ID     =
+EXTRA_SEARCH_MAPPINGS  =
+#---------------------------------------------------------------------------
+# Configuration options related to the LaTeX output
+#---------------------------------------------------------------------------
+GENERATE_LATEX         = NO
+LATEX_OUTPUT           = latex
+LATEX_CMD_NAME         = latex
+MAKEINDEX_CMD_NAME     = makeindex
+LATEX_MAKEINDEX_CMD    = makeindex
+COMPACT_LATEX          = NO
+PAPER_TYPE             = a4wide
+EXTRA_PACKAGES         = float
+LATEX_HEADER           =
+LATEX_FOOTER           =
+LATEX_EXTRA_STYLESHEET =
+LATEX_EXTRA_FILES      =
+PDF_HYPERLINKS         = YES
+USE_PDFLATEX           = YES
+LATEX_BATCHMODE        = YES
+LATEX_HIDE_INDICES     = NO
+LATEX_SOURCE_CODE      = YES
+LATEX_BIB_STYLE        = plain
+LATEX_TIMESTAMP        = NO
+LATEX_EMOJI_DIRECTORY  =
+#---------------------------------------------------------------------------
+# Configuration options related to the RTF output
+#---------------------------------------------------------------------------
+GENERATE_RTF           = NO
+RTF_OUTPUT             = rtf
+COMPACT_RTF            = NO
+RTF_HYPERLINKS         = NO
+RTF_STYLESHEET_FILE    =
+RTF_EXTENSIONS_FILE    =
+RTF_SOURCE_CODE        = NO
+#---------------------------------------------------------------------------
+# Configuration options related to the man page output
+#---------------------------------------------------------------------------
+GENERATE_MAN           = NO
+MAN_OUTPUT             = man
+MAN_EXTENSION          = .3
+MAN_SUBDIR             =
+MAN_LINKS              = YES
+#---------------------------------------------------------------------------
+# Configuration options related to the XML output
+#---------------------------------------------------------------------------
+GENERATE_XML           = YES
+XML_OUTPUT             = xml
+XML_PROGRAMLISTING     = YES
+XML_NS_MEMB_FILE_SCOPE = YES
+#---------------------------------------------------------------------------
+# Configuration options related to the DOCBOOK output
+#---------------------------------------------------------------------------
+GENERATE_DOCBOOK       = NO
+DOCBOOK_OUTPUT         = docbook
+DOCBOOK_PROGRAMLISTING = NO
+#---------------------------------------------------------------------------
+# Configuration options for the AutoGen Definitions output
+#---------------------------------------------------------------------------
+GENERATE_AUTOGEN_DEF   = NO
+#---------------------------------------------------------------------------
+# Configuration options related to the Perl module output
+#---------------------------------------------------------------------------
+GENERATE_PERLMOD       = NO
+PERLMOD_LATEX          = NO
+PERLMOD_PRETTY         = YES
+PERLMOD_MAKEVAR_PREFIX =
+#---------------------------------------------------------------------------
+# Configuration options related to the preprocessor
+#---------------------------------------------------------------------------
+ENABLE_PREPROCESSING   = YES
+MACRO_EXPANSION        = YES
+EXPAND_ONLY_PREDEF     = NO
+SEARCH_INCLUDES        = YES
+INCLUDE_PATH           = ../../source/lib/rocprof-sys-user
+INCLUDE_FILE_PATTERNS  = *.h \
+                         *.hpp
+PREDEFINED             = ROCPROFSYS_PUBLIC_API= \
+                         ROCPROFSYS_HIDDEN_API= \
+                         "ROCPROFSYS_ATTRIBUTE(...)=" \
+                         "ROCPROFSYS_VISIBILITY(...)=" \
+                         "__attribute__(x)=" \
+                         "__declspec(x)=" \
+                         "size_t=unsigned long" \
+                         "uintptr_t=unsigned long" \
+                         DOXYGEN_SHOULD_SKIP_THIS
+EXPAND_AS_DEFINED      =
+SKIP_FUNCTION_MACROS   = NO
+#---------------------------------------------------------------------------
+# Configuration options related to external references
+#---------------------------------------------------------------------------
+TAGFILES               =
+GENERATE_TAGFILE       = html/tagfile.xml
+ALLEXTERNALS           = NO
+EXTERNAL_GROUPS        = YES
+EXTERNAL_PAGES         = YES
+#---------------------------------------------------------------------------
+# Configuration options related to the dot tool
+#---------------------------------------------------------------------------
+CLASS_DIAGRAMS         = YES
+DIA_PATH               =
+HIDE_UNDOC_RELATIONS   = NO
+HAVE_DOT               = NO
+DOT_NUM_THREADS        = 0
+DOT_FONTNAME           = Helvetica
+DOT_FONTSIZE           = 12
+DOT_FONTPATH           =
+CLASS_GRAPH            = NO
+COLLABORATION_GRAPH    = YES
+GROUP_GRAPHS           = YES
+UML_LOOK               = YES
+UML_LIMIT_NUM_FIELDS   = 10
+TEMPLATE_RELATIONS     = YES
+INCLUDE_GRAPH          = YES
+INCLUDED_BY_GRAPH      = YES
+CALL_GRAPH             = NO
+CALLER_GRAPH           = NO
+GRAPHICAL_HIERARCHY    = YES
+DIRECTORY_GRAPH        = YES
+DOT_IMAGE_FORMAT       = svg
+INTERACTIVE_SVG        = YES
+DOT_PATH               = /usr/bin/dot
+DOTFILE_DIRS           =
+MSCFILE_DIRS           =
+DIAFILE_DIRS           =
+PLANTUML_JAR_PATH      =
+PLANTUML_CFG_FILE      =
+PLANTUML_INCLUDE_PATH  =
+DOT_GRAPH_MAX_NODES    = 50
+MAX_DOT_GRAPH_DEPTH    = 0
+DOT_TRANSPARENT        = NO
+DOT_MULTI_TARGETS      = YES
+GENERATE_LEGEND        = YES
+DOT_CLEANUP            = YES

projects/rocprofiler-systems/docs/how-to/configuring-validating-environment.rst

@@ -0,0 +1,71 @@
+.. meta::
+   :description: ROCm Systems Profiler environment validation documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, environment, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Configuring and validating the environment
+****************************************************
+
+After installing `ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_, additional steps are required to set up
+and validate the environment.
+
+.. note::
+
+   The following instructions use the installation path ``/opt/rocprofiler-systems``. If
+   ROCm Systems Profiler is installed elsewhere, substitute the actual installation path.
+
+Configuring the environment
+========================================
+
+After ROCm Systems Profiler is installed, source the ``setup-env.sh`` script to prefix the
+``PATH``, ``LD_LIBRARY_PATH``, and other environment variables:
+
+.. code-block:: shell
+
+   source /opt/rocprofiler-systems/share/rocprofiler-systems/setup-env.sh
+
+Alternatively, if environment modules are supported, add the ``<prefix>/share/modulefiles`` directory
+to ``MODULEPATH``:
+
+.. code-block:: shell
+
+   module use /opt/rocprofiler-systems/share/modulefiles
+
+.. note::
+
+   As an alternative, the above line can be added to the ``${HOME}/.modulerc`` file.
+
+After ROCm Systems Profiler has been added to the ``MODULEPATH``, it can be loaded
+using ``module load rocprofiler-systems/<VERSION>`` and unloaded using ``module unload rocprofiler-systems/<VERSION>``.
+
+.. code-block:: shell
+
+   module load rocprofiler-systems/1.0.0
+   module unload rocprofiler-systems/1.0.0
+
+.. note::
+
+   You might also need to add the path to the ROCm libraries to ``LD_LIBRARY_PATH``,
+   for example, ``export LD_LIBRARY_PATH=/opt/rocm/lib:${LD_LIBRARY_PATH}``
+
+Validating the environment configuration
+========================================
+
+If the following commands all run successfully with the expected output,
+then you are ready to use ROCm Systems Profiler:
+
+.. code-block:: shell
+
+   which rocprof-sys
+   which rocprof-sys-avail
+   which rocprof-sys-sample
+   rocprof-sys-instrument --help
+   rocprof-sys-avail --all
+   rocprof-sys-sample --help
+
+If ROCm Systems Profiler was built with Python support, validate these additional commands:
+
+.. code-block:: shell
+
+   which rocprof-sys-python
+   rocprof-sys-python --help

projects/rocprofiler-systems/docs/how-to/general-tips-using-rocprof-sys.rst

@@ -0,0 +1,60 @@
+.. meta::
+   :description: ROCm Systems Profiler general tips and usage documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, tips, how to, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+********************************************
+General tips for using ROCm Systems Profiler
+********************************************
+
+Follow these general guidelines when using ROCm Systems Profiler. For an explanation of the terms used in this topic, see
+the :doc:`ROCm Systems Profiler glossary <../reference/rocprof-sys-glossary>`.
+
+* Use ``rocprof-sys-avail`` to look up configuration settings, hardware counters, and data collection components
+
+  * Use the ``-d`` flag for descriptions
+
+* Generate a default configuration with ``rocprof-sys-avail -G ${HOME}/.rocprof-sys.cfg`` and adjust it
+  to the desired default behavior
+* **Decide whether binary instrumentation, statistical sampling, or both** provides the desired performance data (for non-Python applications)
+* Compile code with optimization enabled (``-O2`` or higher), disable asserts (i.e. ``-DNDEBUG``), and include debug info (for instance, ``-g1`` at a minimum)
+
+  * Compiling with debug info does not slow down the code, it only increases compile time and the size of the binary
+  * In CMake, this is generally done with the settings ``CMAKE_BUILD_TYPE=RelWithDebInfo`` or ``CMAKE_BUILD_TYPE=Release`` and ``CMAKE_<LANG>_FLAGS=-g1``
+
+* **Use binary instrumentation for characterizing the performance of every invocation of specific functions**
+* **Use statistical sampling to characterize the performance of the entire application while minimizing overhead**
+* Enable statistical sampling after binary instrumentation to help "fill in the gaps" between instrumented regions
+* Use the user API to create custom regions and enable/disable ROCm Systems Profiler for specific processes, threads, and regions
+* Dynamic symbol interception, callback APIs, and the user API are always available with binary instrumentation and sampling
+
+  * Dynamic symbol interception and callback APIs are (generally) controlled through ``ROCPROFSYS_USE_<API>``
+    options, for example, ``ROCPROFSYS_USE_KOKKOSP`` and ``ROCPROFSYS_USE_OMPT`` enable Kokkos-Tools and OpenMP-Tools
+    callbacks, respectively
+
+* When generically seeking regions for performance improvement:
+
+  * **Start off by collecting a flat profile**
+  * Look for functions with high call counts, large cumulative runtimes/values, or large standard deviations
+
+    * When call counts are high, improving the performance of this function or "inlining" the function can result in quick and easy performance improvements
+    * When the standard deviation is high, collect a hierarchical profile and see if the high variation can be attributable to the calling context.
+      In this scenario, consider creating a specialized version of the function for the longer-running contexts
+
+  * **Collect a hierarchical profile** and verify the functions that are part of the "critical path" of your
+    application, as indicated in the flat profile
+
+    * For example, functions with high call counts but which are part of a "setup" or "post-processing"
+      phase that does not consume much time relative to the overall time are generally a lower priority for optimization
+
+* **Use the information from the profiles when analyzing detailed traces**
+* When using binary instrumentation in "trace" mode, **binary rewrites are preferable to runtime instrumentation**.
+
+  * Binary rewrites only instrument the functions defined in the target binary, whereas runtime instrumentation might instrument functions defined in the shared libraries which are linked into the target binary
+
+* When using binary instrumentation with MPI, avoid runtime instrumentation
+
+  * Runtime instrumentation requires a fork and a ``ptrace``, which is generally incompatible with how MPI applications spawn processes
+  * Perform a binary rewrite of the executable (and optionally, libraries used by the executable) using MPI and run
+    the generated instrumented executable using ``rocprof-sys-run`` instead of the original.
+    For example, instead of ``mpirun -n 2 ./myexe``, use ``mpirun -n 2 rocprof-sys-run -- ./myexe.inst``, where
+    ``myexe.inst`` is the instrumented ``myexe`` executable that was generated.

projects/rocprofiler-systems/docs/how-to/instrumenting-rewriting-binary-application.rst

@@ -0,0 +1,939 @@
+.. meta::
+   :description: ROCm Systems Profiler binary instrumentation and rewrite documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, binary instrumentation, binary rewrite, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Instrumenting and rewriting a binary application
+****************************************************
+
+There are three ways to perform instrumentation with the ``rocprof-sys-instrument`` executable:
+
+* Runtime instrumentation
+* Attaching to an already running process
+* Binary rewrite
+
+Here is a comparison of the three modes:
+
+* Runtime instrumentation of the application using the ``rocprof-sys-instrument`` executable
+  (analogous to ``gdb --args <program> <args>``)
+
+  * This mode is the default if neither the ``-p`` nor ``-o`` command-line options are used
+  * Runtime instrumentation supports instrumenting not only the target executable but also
+    the shared libraries loaded by the target executable. Consequently, this mode consumes more memory,
+    takes longer to perform the instrumentation, and tends to add more significant overhead to the
+    runtime of the application.
+  * This mode is recommended if you want to analyze not only the performance of your executable and/or
+    libraries but also the performance of the library dependencies
+
+* Attaching to a process that is currently running (analogous to ``gdb -p <PID>``)
+
+  * This mode is activated using ``-p <PID>``
+  * The same caveats from the first example apply with respect to memory and overhead
+
+  .. note::
+
+     Attaching to a running process is an alpha feature and detaching from the target process
+     without ending the target process is not currently supported.
+
+* Binary rewrite to generate a new executable or library with the instrumentation built-in
+
+  * This mode is activated through the ``-o <output-file>`` option
+  * Binary rewriting is limited to the text section of the target executable or library. It does not instrument
+    the dynamically-linked libraries. Consequently, this mode performs the
+    instrumentation significantly faster
+    and has a much lower overhead when running the instrumented executable and libraries.
+  * Binary rewriting is the recommended mode when the target executable uses
+    process-level parallelism (for example, MPI)
+  * If the target executable has a minimal ``main`` routine and the bulk of your
+    application is in one specific dynamic library,
+    see :ref:`binary-rewriting-library-label` for help
+
+The rocprof-sys-instrument executable
+========================================
+
+Instrumentation is performed with the ``rocprof-sys-instrument`` executable. For more details, use the ``-h`` or ``--help`` option to
+view the help menu.
+
+.. code-block:: shell
+
+   $ rocprof-sys-instrument --help
+   [rocprof-sys-instrument] Usage: rocprof-sys-instrument [ --help (count: 0, dtype: bool)
+                                                      --version (count: 0, dtype: bool)
+                                                      --verbose (max: 1, dtype: bool)
+                                                      --error (max: 1, dtype: boolean)
+                                                      --debug (max: 1, dtype: bool)
+                                                      --log (count: 1)
+                                                      --log-file (count: 1)
+                                                      --simulate (max: 1, dtype: boolean)
+                                                      --print-format (min: 1, dtype: string)
+                                                      --print-dir (count: 1, dtype: string)
+                                                      --print-available (count: 1)
+                                                      --print-instrumented (count: 1)
+                                                      --print-coverage (count: 1)
+                                                      --print-excluded (count: 1)
+                                                      --print-overlapping (count: 1)
+                                                      --print-instructions (max: 1, dtype: bool)
+                                                      --output (min: 0, dtype: string)
+                                                      --pid (count: 1, dtype: int)
+                                                      --mode (count: 1)
+                                                      --force (max: 1, dtype: bool)
+                                                      --command (count: 1)
+                                                      --prefer (count: 1)
+                                                      --library (count: unlimited)
+                                                      --main-function (count: 1)
+                                                      --load (count: unlimited, dtype: string)
+                                                      --load-instr (count: unlimited, dtype: filepath)
+                                                      --init-functions (count: unlimited, dtype: string)
+                                                      --fini-functions (count: unlimited, dtype: string)
+                                                      --all-functions (max: 1, dtype: boolean)
+                                                      --function-include (count: unlimited)
+                                                      --function-exclude (count: unlimited)
+                                                      --function-restrict (count: unlimited)
+                                                      --caller-include (count: unlimited)
+                                                      --module-include (count: unlimited)
+                                                      --module-exclude (count: unlimited)
+                                                      --module-restrict (count: unlimited)
+                                                      --internal-function-include (count: unlimited)
+                                                      --internal-module-include (count: unlimited)
+                                                      --instruction-exclude (count: unlimited)
+                                                      --internal-library-deps (min: 0, dtype: boolean)
+                                                      --internal-library-append (count: unlimited)
+                                                      --internal-library-remove (count: unlimited)
+                                                      --linkage (min: 1)
+                                                      --visibility (min: 1)
+                                                      --label (count: unlimited, dtype: string)
+                                                      --config (min: 1, dtype: string)
+                                                      --default-components (count: unlimited, dtype: string)
+                                                      --env (count: unlimited)
+                                                      --mpi (max: 1, dtype: bool)
+                                                      --instrument-loops (max: 1, dtype: boolean)
+                                                      --min-instructions (count: 1, dtype: int)
+                                                      --min-address-range (count: 1, dtype: int)
+                                                      --min-instructions-loop (count: 1, dtype: int)
+                                                      --min-address-range-loop (count: 1, dtype: int)
+                                                      --coverage (max: 1, dtype: bool)
+                                                      --dynamic-callsites (max: 1, dtype: boolean)
+                                                      --traps (max: 1, dtype: boolean)
+                                                      --loop-traps (max: 1, dtype: boolean)
+                                                      --allow-overlapping (max: 1, dtype: bool)
+                                                      --parse-all-modules (max: 1, dtype: bool)
+                                                      --batch-size (count: 1, dtype: int)
+                                                      --dyninst-rt (min: 1, dtype: filepath)
+                                                      --dyninst-options (count: unlimited)
+                                                      ] -- <CMD> <ARGS>
+
+   Options:
+      -h, -?, --help                 Shows this page
+      --version                      Prints the version and exit
+
+      [DEBUG OPTIONS]
+
+      -v, --verbose                  Verbose output
+      -e, --error                    All warnings produce runtime errors
+      --debug                        Debug output
+      --log                          Number of log entries to display after an error. Any value < 0 will emit the entire log
+      --log-file                     Write the log out the specified file during the run
+      --simulate                     Exit after outputting diagnostic {available,instrumented,excluded,overlapping} module
+                                    function lists, e.g. available.txt
+      --print-format [ json | txt | xml ]
+                                    Output format for diagnostic {available,instrumented,excluded,overlapping} module
+                                    function lists, e.g. {print-dir}/available.txt
+      --print-dir                    Output directory for diagnostic {available,instrumented,excluded,overlapping} module
+                                    function lists, e.g. {print-dir}/available.txt
+      --print-available [ functions | functions+ | modules | pair | pair+ ]
+                                    Print the available entities for instrumentation (functions, modules, or module-function
+                                    pair) to stdout after applying regular expressions
+      --print-instrumented [ functions | functions+ | modules | pair | pair+ ]
+                                    Print the instrumented entities (functions, modules, or module-function pair) to stdout
+                                    after applying regular expressions
+      --print-coverage [ functions | functions+ | modules | pair | pair+ ]
+                                    Print the instrumented coverage entities (functions, modules, or module-function pair) to
+                                    stdout after applying regular expressions
+      --print-excluded [ functions | functions+ | modules | pair | pair+ ]
+                                    Print the entities for instrumentation (functions, modules, or module-function pair)
+                                    which are excluded from the instrumentation to stdout after applying regular expressions
+      --print-overlapping [ functions | functions+ | modules | pair | pair+ ]
+                                    Print the entities for instrumentation (functions, modules, or module-function pair)
+                                    which overlap other function calls or have multiple entry points to stdout after applying
+                                    regular expressions
+      --print-instructions           Print the instructions for each basic-block in the JSON/XML outputs
+
+      [MODE OPTIONS]
+
+      -o, --output                   Enable generation of a new executable (binary-rewrite). If a filename is not provided,
+                                    rocprof-sys will use the basename and output to the cwd, unless the target binary is in the
+                                    cwd. In the latter case, rocprof-sys will either use ${PWD}/<basename>.inst (non-libraries)
+                                    or ${PWD}/instrumented/<basename> (libraries)
+      -p, --pid                      Connect to running process
+      -M, --mode [ coverage | sampling | trace ]
+                                    Instrumentation mode. \'trace\' mode instruments the selected functions, \'sampling\' mode
+                                    only instruments the main function to start and stop the sampler.
+      -f, --force                    Force the command-line argument configuration, i.e. don't get cute. Useful for forcing
+                                    runtime instrumentation of an executable that [A] Dyninst thinks is a library after
+                                    reading ELF and [B] whose name makes it look like a library (e.g. starts with 'lib'
+                                    and/or ends in \'.so\', \'.so.*\', or \'.a\')
+      -c, --command                  Input executable and arguments (if \'-- <CMD>\' not provided)
+
+      [LIBRARY OPTIONS]
+
+      --prefer [ shared | static ]   Prefer this library types when available
+      -L, --library                  Libraries with instrumentation routines (default: "librocprof-sys-dl")
+      -m, --main-function            The primary function to instrument around, e.g. \'main\'
+      --load                         Supplemental instrumentation library names w/o extension (e.g. \'libinstr\' for
+                                    \'libinstr.so\' or \'libinstr.a\')
+      --load-instr                   Load {available,instrumented,excluded,overlapping}-instr JSON or XML file(s) and override
+                                    what is read from the binary
+      --init-functions               Initialization function(s) for supplemental instrumentation libraries (see \'--load\'
+                                    option)
+      --fini-functions               Finalization function(s) for supplemental instrumentation libraries (see \'--load\' option)
+      --all-functions                When finding functions, include the functions which are not instrumentable. This is
+                                    purely diagnostic for the available/excluded functions output
+
+      [SYMBOL SELECTION OPTIONS]
+
+      -I, --function-include         Regex(es) for including functions (despite heuristics)
+      -E, --function-exclude         Regex(es) for excluding functions (always applied)
+      -R, --function-restrict        Regex(es) for restricting functions only to those that match the provided
+                                    regular-expressions
+      --caller-include               Regex(es) for including functions that call the listed functions (despite heuristics)
+      -MI, --module-include          Regex(es) for selecting modules/files/libraries (despite heuristics)
+      -ME, --module-exclude          Regex(es) for excluding modules/files/libraries (always applied)
+      -MR, --module-restrict         Regex(es) for restricting modules/files/libraries only to those that match the provided
+                                    regular-expressions
+      --internal-function-include    Regex(es) for including functions which are (likely) utilized by rocprof-sys itself. Use
+                                    this option with care.
+      --internal-module-include      Regex(es) for including modules/libraries which are (likely) utilized by rocprof-sys
+                                    itself. Use this option with care.
+      --instruction-exclude          Regex(es) for excluding functions containing certain instructions
+      --internal-library-deps        Treat the libraries linked to the internal libraries as internal libraries. This increase
+                                    the internal library processing time and consume more memory (so use with care) but may
+                                    be useful when the application uses Boost libraries and Dyninst is dynamically linked
+                                    against the same boost libraries
+      --internal-library-append      Append to the list of libraries which rocprof-sys treats as being used internally, e.g.
+                                    ROCm Systems Profiler will find all the symbols in this library and prevent them from being
+                                    instrumented.
+      --internal-library-remove [ ld-linux-x86-64.so.2
+                                 libBrokenLocale.so.1
+                                 libanl.so.1
+                                 libbfd.so
+                                 libbz2.so
+                                 libc.so.6
+                                 libcaliper.so
+                                 libcommon.so
+                                 libcrypt.so.1
+                                 libdl.so.2
+                                 libdw.so
+                                 libdwarf.so
+                                 libdyninstAPI_RT.so
+                                 libelf.so
+                                 libgcc_s.so.1
+                                 libgotcha.so
+                                 liblikwid.so
+                                 liblzma.so
+                                 libnsl.so.1
+                                 libnss_compat.so.2
+                                 libnss_db.so.2
+                                 libnss_dns.so.2
+                                 libnss_files.so.2
+                                 libnss_hesiod.so.2
+                                 libnss_ldap.so.2
+                                 libnss_nis.so.2
+                                 libnss_nisplus.so.2
+                                 libnss_test1.so.2
+                                 libnss_test2.so.2
+                                 libpapi.so
+                                 libpfm.so
+                                 libprofiler.so
+                                 libpthread.so.0
+                                 libresolv.so.2
+                                 libamd_smi64.so
+                                 librocprofiler-sdk.so
+                                 librt.so.1
+                                 libstdc++.so.6
+                                 libtbb.so
+                                 libtbbmalloc.so
+                                 libtbbmalloc_proxy.so
+                                 libtcmalloc.so
+                                 libtcmalloc_and_profiler.so
+                                 libtcmalloc_debug.so
+                                 libtcmalloc_minimal.so
+                                 libtcmalloc_minimal_debug.so
+                                 libthread_db.so.1
+                                 libunwind-coredump.so
+                                 libunwind-generic.so
+                                 libunwind-ptrace.so
+                                 libunwind-setjmp.so
+                                 libunwind-x86_64.so
+                                 libunwind.so
+                                 libutil.so.1
+                                 libz.so
+                                 libzstd.so ]
+                                    Remove the specified libraries from being treated as being used internally, e.g.
+                                    ROCm System Profiler will permit all the symbols in these libraries to be eligible for
+                                    instrumentation.
+      --linkage [ global | local | unique | unknown | weak ]
+                                    Only instrument functions with specified linkage (default: global, local, unique)
+      --visibility [ default | hidden | internal | protected | unknown ]
+                                    Only instrument functions with specified visibility (default: default, internal, hidden,
+                                    protected)
+
+      [RUNTIME OPTIONS]
+
+      --label [ args | file | line | return ]
+                                    Labeling info for functions. By default, just the function name is recorded. Use these
+                                    options to gain more information about the function signature or location of the
+                                    functions
+      -C, --config                   Read in a configuration file and encode these values as the defaults in the executable
+      -d, --default-components       Default components to instrument (only useful when timemory is enabled in rocprof-sys
+                                    library)
+      --env                          Environment variables to add to the runtime in form VARIABLE=VALUE. E.g. use \'--env
+                                    ROCPROFSYS_PROFILE=ON\' to default to using timemory instead of perfetto
+      --mpi                          Enable MPI support (requires rocprof-sys built w/ full or partial MPI support). NOTE: this
+                                    will automatically be activated if MPI_Init, MPI_Init_thread, MPI_Finalize,
+                                    MPI_Comm_rank, or MPI_Comm_size are found in the symbol table of target
+
+      [GRANULARITY OPTIONS]
+
+      -l, --instrument-loops         Instrument at the loop level
+      -i, --min-instructions         If the number of instructions in a function is less than this value, exclude it from
+                                    instrumentation
+      -r, --min-address-range        If the address range of a function is less than this value, exclude it from
+                                    instrumentation
+      --min-instructions-loop        If the number of instructions in a function containing a loop is less than this value,
+                                    exclude it from instrumentation
+      --min-address-range-loop       If the address range of a function containing a loop is less than this value, exclude it
+                                    from instrumentation
+      --coverage [ basic_block | function | none ]
+                                    Enable recording the code coverage. If instrumenting in coverage mode (\'-M converage\'),
+                                    this simply specifies the granularity. If instrumenting in trace or sampling mode, this
+                                    enables recording code-coverage in addition to the instrumentation of that mode (if any).
+      --dynamic-callsites            Force instrumentation if a function has dynamic callsites (e.g. function pointers)
+      --traps                        Instrument points which require using a trap. On the x86 architecture, because
+                                    instructions are of variable size, the instruction at a point may be too small for
+                                    Dyninst to replace it with the normal code sequence used to call instrumentation. Also,
+                                    when instrumentation is placed at points other than subroutine entry, exit, or call
+                                    points, traps may be used to ensure the instrumentation fits. In this case, Dyninst
+                                    replaces the instruction with a single-byte instruction that generates a trap.
+      --loop-traps                   Instrument points within a loop which require using a trap (only relevant when
+                                    --instrument-loops is enabled).
+      --allow-overlapping            Allow dyninst to instrument either multiple functions which overlap (share part of same
+                                    function body) or single functions with multiple entry points. For more info, see Section
+                                    2 of the DyninstAPI documentation.
+      --parse-all-modules            By default, rocprof-sys simply requests Dyninst to provide all the procedures in the
+                                    application image. If this option is enabled, rocprof-sys will iterate over all the modules
+                                    and extract the functions. Theoretically, it should be the same but the data is slightly
+                                    different, possibly due to weak binding scopes. In general, enabling option will probably
+                                    have no visible effect
+
+      [DYNINST OPTIONS]
+
+      -b, --batch-size               Dyninst supports batch insertion of multiple points during runtime instrumentation. If
+                                    one large batch insertion fails, this value will be used to create smaller batches.
+                                    Larger batches generally decrease the instrumentation time
+      --dyninst-rt                   Path(s) to the dyninstAPI_RT library
+      --dyninst-options [ BaseTrampDeletion
+                           DebugParsing
+                           DelayedParsing
+                           InstrStackFrames
+                           MergeTramp
+                           SaveFPR
+                           TrampRecursive
+                           TypeChecking ]
+      Advanced dyninst options: BPatch::set<OPTION>(bool), e.g. bpatch->setTrampRecursive(true)
+
+``rocprof-sys-instrument`` uses a similar syntax as LLVM to separate command-line arguments from the
+application's arguments. It uses a standalone
+double-hyphen (``--``) as a separator.
+All arguments preceding the double-hyphen
+are interpreted as belonging to ROCm Systems Profiler and all arguments following the
+double-hyphen are interpreted as being part of the
+application and its arguments. In binary rewrite mode, all application arguments after the first argument
+are ignored. As an example, ``./rocprof-sys-instrument -o ls.inst -- ls -l`` interprets ``ls`` as
+the target to instrument, ignoring the ``-l`` argument,
+and generates a ``ls.inst`` executable that you can subsequently run using the
+``rocprof-sys-run -- ls.inst -l`` command.
+
+Runtime instrumentation example
+========================================
+
+The following example shows how to enable runtime instrumentation.
+
+.. code-block:: shell
+
+   rocprof-sys-instrument <rocprof-sys-options> -- <exe> [<exe-options>...]
+
+Attaching to a running process
+========================================
+
+Use the following command to attach to an active process.
+
+.. code-block:: shell
+
+   rocprof-sys-instrument <rocprof-sys-options> -p <PID> -- <exe-name>
+
+Binary rewrite
+========================================
+
+This example demonstrates how to rewrite a binary.
+
+.. code-block:: shell
+
+   rocprof-sys-instrument <rocprof-sys-options> -o <name-of-new-exe-or-library> -- <exe-or-library>
+
+.. _binary-rewriting-library-label:
+
+Binary rewrite of a library
+-----------------------------------
+
+Many applications bundle the bulk of their functionality into one or more
+dynamic libraries and have a relatively simple ``main``
+which links to these libraries and serves as the "driver" for
+setting up the workflow. If you perform a binary rewrite of an
+executable like this and find there is insufficient information, you
+can either switch to runtime instrumentation or perform a
+binary rewrite on the relevant libraries.
+
+Support for stand-alone binary rewriting of a dynamic library without a binary rewrite of
+the executable is a beta feature.
+In general, it is supported as long as the library contains the ``_init`` and
+``_fini`` symbols but these symbols are not
+standardized to the extent of ``main`` in an executable.
+
+Here is the recommended workflow for the binary rewrite of a library:
+
+#. Determine the names of the dynamically linked libraries of interest using ``ldd``
+#. Generate a binary rewrite of the executable
+#. Generate a binary rewrite of the desired libraries with the same base name as the
+   original library, for example, ``libfoo.so.2`` instead of ``libfoo.so``,  and output the instrumented
+   library into a different folder than the original library.
+
+#. Prefix the ``LD_LIBRARY_PATH`` executable with the output folder from the previous step
+#. Use ``ldd`` to verify that the instrumented executable can resolve the location of the instrumented library
+
+Binary rewrite of a library example
+-----------------------------------
+
+The ``foo`` executable is dynamically linked to ``libfoo.so.2``:
+
+.. code-block:: shell
+
+   $ pwd
+   /home/user
+   $ which foo
+   /usr/local/bin/foo
+   $ ldd /usr/local/bin/foo
+         ...
+         libfoo.so.2 => /usr/local/lib/libfoo.so.2 (...)
+         ...
+
+Generate binary rewrites of ``foo`` and ``libfoo.so.2``:
+
+.. code-block:: shell
+
+   rocprof-sys-instrument -o ./foo.inst -- foo
+   rocprof-sys-instrument -o ./libfoo.so.2 -- /usr/local/lib/libfoo.so.2
+
+At this point, the instrumented ``foo.inst`` executable still dynamically loads the
+original ``libfoo.so.2`` in ``/usr/local/lib``:
+
+.. code-block:: shell
+
+   $ ldd ./foo.inst
+         ...
+         libfoo.so.2 => /usr/local/lib/libfoo.so.2 (...)
+         ...
+
+Prefix the ``LD_LIBRARY_PATH`` environment variable with the folder containing
+the instrumented ``libfoo.so.2``:
+
+.. code-block:: shell
+
+   export LD_LIBRARY_PATH=/home/user:${LD_LIBRARY_PATH}
+
+``foo.inst`` now loads the instrumented library when it runs:
+
+.. code-block:: shell
+
+   $ ldd ./foo.inst
+         ...
+         libfoo.so.2 => /home/user/libfoo.so.2 (...)
+         ...
+
+Selective instrumentation
+========================================
+
+The default behavior of ``rocprof-sys-instrument`` does not instrument every symbol in the binary.
+The default rules are:
+
+* Skip instrumenting dynamic call-sites (such as function pointers)
+
+  * The ``--dynamic-callsites`` option forces instrumentation for all dynamic call-sites
+
+* The cost of a function can be loosely approximated by the number of
+  instructions. By default, ``rocprof-sys-instrument`` only instruments functions
+  with at least 1024 instructions
+
+  * The  ``--min-instructions`` option modifies this heuristic for all functions which do not contain loops
+  * The ``--min-instructions-loop`` option modifies this heuristic for functions which contain loops.
+
+* The cost of a function can be also be loosely approximated by the size of the function
+  in the binary so this heuristic can be used in lieu of or in addition to the
+  minimum number of instructions
+
+  * The ``--min-address-range`` option modifies this heuristic for all functions which do not contain loops
+  * The ``--min-address-range-loop`` option modifies this heuristic for functions which contain loops
+
+* Skip instrumentation points which require using a trap
+
+  * See the description for the ``--traps`` and ``--loop-traps`` options for more information
+
+* Skip instrumenting loops within the body of a function
+
+  * The ``--instrument-loops`` option enables this behavior
+
+* Skip instrumenting functions with overlapping function bodies and single
+  functions with multiple entry point
+
+  * These behaviors arise from various optimizations. Enable instrumenting for these functions
+    by using the ``--allow-overlapping`` option
+
+.. note::
+
+   The separate loop options ``--min-instructions-loop`` and ``--min-address-range-loop``
+   are provided because functions with loops can be compact in the binary while also being costly
+
+Viewing the available, instrumented, excluded, and overlapping functions
+-------------------------------------------------------------------------
+
+Whenever ``rocprof-sys-instrument`` runs with a verbosity of zero or higher,
+it generates files that detail which functions
+were available for instrumentation (along with the module they were defined in), actually instrumented,
+excluded, and which contained overlapping function bodies.
+By default, these files are saved to the ``rocprof-sys-<NAME>-output`` folder
+where ``<NAME>`` is the base name of the targeted binary (or
+the base name of the resulting executable in the case of binary rewrite). For example,
+``rocprof-sys-instrument -- ls`` outputs these files to ``rocprof-sys-ls-output``
+whereas ``rocprof-sys-instrument -o ls.inst -- ls`` places them in ``rocprof-sys-ls.inst-output``.
+
+To generate these files without running or generating an
+executable, use the ``--simulate`` option:
+
+.. code-block:: shell
+
+   rocprof-sys-instrument --simulate -- foo
+   rocprof-sys-instrument --simulate -o foo.inst -- foo
+
+Excluding and including modules and functions
+----------------------------------------------
+
+ROCm Systems Profiler has a set of six command-line options which each accept one or more
+regular expressions for customizing the scope of which module and/or functions are
+instrumented. Multiple regex patterns per option are treated as an OR operation,
+for example, ``--module-include libfoo libbar`` is effectively the same as ``--module-include 'libfoo|libbar'``.
+
+To force the inclusion of certain modules and/or function
+without changing any of the heuristics, use the ``--module-include`` and/or ``--function-include`` options.
+These options do not exclude modules or functions which do
+not satisfy their regular expression.
+
+To narrow the scope of the instrumentation to a specific set
+of libraries and/or functions, use the ``--module-restrict`` and ``--function-restrict`` options.
+These options let you exclusively select the union of one or more
+regular expressions, regardless of whether or not the functions satisfy the
+previously-mentioned default heuristics. Any function or module that is not within
+the union of these regular expressions is excluded from instrumentation.
+
+To avoid instrumenting a set of modules and/or functions,
+use the ``--module-exclude`` and ``--function-exclude`` options.
+These options are always applied, even if the module or function
+satisfies the "restrict" or "include" regular expression.
+
+.. _available-module-function-output:
+
+An example of the available module and function info output
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+.. code-block:: shell
+
+   rocprof-sys-instrument -o lulesh.inst --label file line args --simulate -- lulesh
+
+.. code-block:: shell
+
+   AddressRange  Module                                    Function                                                                                 FunctionSignature
+           9165  ../examples/lulesh/lulesh-comm.cc         CommMonoQ                                                                                CommMonoQ(domain) [lulesh-comm.cc:1891]
+           3396  ../examples/lulesh/lulesh-comm.cc         CommRecv                                                                                 CommRecv(domain, int, Index_t, Index_t, Index_t, Index_t, bool, bool) [lulesh...
+           8666  ../examples/lulesh/lulesh-comm.cc         CommSBN                                                                                  CommSBN(domain, int, Domain_member *) [lulesh-comm.cc:926]
+          10212  ../examples/lulesh/lulesh-comm.cc         CommSend                                                                                 CommSend(domain, int, Index_t, Domain_member *, Index_t, Index_t, Index_t, bo...
+           6823  ../examples/lulesh/lulesh-comm.cc         CommSyncPosVel                                                                           CommSyncPosVel(domain) [lulesh-comm.cc:1404]
+            126  ../examples/lulesh/lulesh-comm.cc         _GLOBAL__sub_I_lulesh_comm.cc                                                            _GLOBAL__sub_I_lulesh_comm.cc() [lulesh-comm.cc]
+            308  ../examples/lulesh/lulesh-init.cc         .omp_outlined..26                                                                        .omp_outlined..26(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            628  ../examples/lulesh/lulesh-init.cc         .omp_outlined..34                                                                        .omp_outlined..34(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            656  ../examples/lulesh/lulesh-init.cc         .omp_outlined..41                                                                        .omp_outlined..41(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            662  ../examples/lulesh/lulesh-init.cc         .omp_outlined..45                                                                        .omp_outlined..45(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            550  ../examples/lulesh/lulesh-init.cc         .omp_outlined..55                                                                        .omp_outlined..55(const , const , const ParallelFor<Kokkos::Impl::ViewFill<Ko...
+            556  ../examples/lulesh/lulesh-init.cc         .omp_outlined..57                                                                        .omp_outlined..57(const , const , const ParallelFor<Kokkos::Impl::ViewFill<Ko...
+            550  ../examples/lulesh/lulesh-init.cc         .omp_outlined..78                                                                        .omp_outlined..78(const , const , const ParallelFor<Kokkos::Impl::ViewFill<Ko...
+            640  ../examples/lulesh/lulesh-init.cc         .omp_outlined..84                                                                        .omp_outlined..84(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            646  ../examples/lulesh/lulesh-init.cc         .omp_outlined..88                                                                        .omp_outlined..88(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+           1840  ../examples/lulesh/lulesh-init.cc         Domain::AllocateElemPersistent                                                           Domain::AllocateElemPersistent(Domain *, Int_t) [lulesh-init.cc:94]
+           1384  ../examples/lulesh/lulesh-init.cc         Domain::AllocateNodePersistent                                                           Domain::AllocateNodePersistent(Domain *, Int_t) [lulesh-init.cc:94]
+           1264  ../examples/lulesh/lulesh-init.cc         Domain::BuildMesh                                                                        Domain::BuildMesh(Domain *, Int_t, Int_t, Int_t) [lulesh-init.cc:308]
+           2312  ../examples/lulesh/lulesh-init.cc         Domain::CreateRegionIndexSets                                                            Domain::CreateRegionIndexSets(Domain *, Int_t, Int_t) [lulesh-init.cc:409]
+           7109  ../examples/lulesh/lulesh-init.cc         Domain::Domain                                                                           Domain::Domain(Domain *, Int_t, Index_t, Index_t, Index_t, Index_t, int, int,...
+           2458  ../examples/lulesh/lulesh-init.cc         Domain::SetupBoundaryConditions                                                          Domain::SetupBoundaryConditions(Domain *, Int_t) [lulesh-init.cc:409]
+            956  ../examples/lulesh/lulesh-init.cc         Domain::SetupCommBuffers                                                                 Domain::SetupCommBuffers(Domain *, Int_t) [lulesh-init.cc]
+           1456  ../examples/lulesh/lulesh-init.cc         Domain::SetupElementConnectivities                                                       Domain::SetupElementConnectivities(Domain *, Int_t) [lulesh-init.cc:409]
+            721  ../examples/lulesh/lulesh-init.cc         Domain::SetupSymmetryPlanes                                                              Domain::SetupSymmetryPlanes(Domain *, Int_t) [lulesh-init.cc:409]
+           1591  ../examples/lulesh/lulesh-init.cc         Domain::SetupThreadSupportStructures                                                     Domain::SetupThreadSupportStructures(Domain *) [lulesh-init.cc:376]
+           1644  ../examples/lulesh/lulesh-init.cc         Domain::~Domain                                                                          Domain::~Domain(Domain *) [lulesh-init.cc:286]
+            218  ../examples/lulesh/lulesh-init.cc         InitMeshDecomp                                                                           InitMeshDecomp(Int_t, Int_t, Int_t *, Int_t *, Int_t *, Int_t *) [lulesh-init...
+            260  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::CommonSubview<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokk...         Kokkos::Impl::CommonSubview<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokk...
+           1786  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::HostIterateTile<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::R...         Kokkos::Impl::HostIterateTile<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::R...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int**...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int**...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int**...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int**...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int*,...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int*,...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int*,...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<int*,...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewFill<Kokkos::View<doubl...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewFill<Kokkos::View<doubl...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewFill<Kokkos::View<doubl...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewFill<Kokkos::View<doubl...
+            330  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewFill<Kokkos::View<doubl...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewFill<Kokkos::View<doubl...
+            522  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelFor<Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::...         Kokkos::Impl::ParallelFor<Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::...
+            232  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ParallelFor<Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::...         Kokkos::Impl::ParallelFor<Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::...
+             49  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewVal...         Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewVal...
+           1476  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::Tile_Loop_Type<2, false, int, void, void>::apply<Kokkos::Impl::...         Kokkos::Impl::Tile_Loop_Type<2, false, int, void, void>::apply<Kokkos::Impl::...
+            555  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::LayoutRight, Kokkos::Devic...         Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::LayoutRight, Kokkos::Devic...
+            613  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::LayoutRight, Kokkos::Devic...         Kokkos::Impl::ViewCopy<Kokkos::View<int**, Kokkos::LayoutRight, Kokkos::Devic...
+            603  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCopy<Kokkos::View<int*, Kokkos::LayoutLeft, Kokkos::Device<...         Kokkos::Impl::ViewCopy<Kokkos::View<int*, Kokkos::LayoutLeft, Kokkos::Device<...
+            604  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCopy<Kokkos::View<int*, Kokkos::LayoutLeft, Kokkos::Device<...         Kokkos::Impl::ViewCopy<Kokkos::View<int*, Kokkos::LayoutLeft, Kokkos::Device<...
+            281  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            281  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            281  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            281  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            281  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            524  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewFill<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::Dev...         Kokkos::Impl::ViewFill<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::Dev...
+            525  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewFill<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::Dev...         Kokkos::Impl::ViewFill<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::Dev...
+            524  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewFill<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::Dev...         Kokkos::Impl::ViewFill<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::Dev...
+            583  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewMapping<Kokkos::ViewTraits<int* [8], Kokkos::LayoutRight>, ...         SharedAllocationRecord<void, void> * Kokkos::Impl::ViewMapping<Kokkos::ViewTr...
+            529  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewMapping<Kokkos::ViewTraits<int*, Kokkos::HostSpace>, void>:...         SharedAllocationRecord<void, void> * Kokkos::Impl::ViewMapping<Kokkos::ViewTr...
+            529  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewMapping<Kokkos::ViewTraits<int*>, void>::allocate_shared<st...         SharedAllocationRecord<void, void> * Kokkos::Impl::ViewMapping<Kokkos::ViewTr...
+            203  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewRemap<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokkos::...         Kokkos::Impl::ViewRemap<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokkos::...
+            331  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewRemap<Kokkos::View<int*>, Kokkos::View<int*>, Kokkos::OpenM...         Kokkos::Impl::ViewRemap<Kokkos::View<int*>, Kokkos::View<int*>, Kokkos::OpenM...
+            461  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::ViewValueFunctor<Kokkos::Device<Kokkos::OpenMP, Kokkos::HostSpa...         enable_if_t<std::is_trivial<int>::value && std::is_trivially_copy_assignable<...
+            353  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::contiguous_fill<Kokkos::OpenMP, double*>                                   Kokkos::Impl::contiguous_fill<Kokkos::OpenMP, double*>(exec_space, dst, value...
+            139  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::contiguous_fill<Kokkos::OpenMP, double, Kokkos::LayoutRight, Ko...         Kokkos::Impl::contiguous_fill<Kokkos::OpenMP, double, Kokkos::LayoutRight, Ko...
+            824  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight, Kokkos::D...         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight, Kokkos::D...
+            824  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight, Kokkos::D...         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight, Kokkos::D...
+            824  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokkos::...         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokkos::...
+            824  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokkos::...         Kokkos::Impl::view_copy<Kokkos::View<int* [8], Kokkos::LayoutRight>, Kokkos::...
+            697  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::view_copy<Kokkos::View<int*, Kokkos::LayoutRight, Kokkos::Devic...         Kokkos::Impl::view_copy<Kokkos::View<int*, Kokkos::LayoutRight, Kokkos::Devic...
+            697  ../examples/lulesh/lulesh-init.cc         Kokkos::Impl::view_copy<Kokkos::View<int*>, Kokkos::View<int*> >                         Kokkos::Impl::view_copy<Kokkos::View<int*>, Kokkos::View<int*> >(dst, src) [l...
+           2036  ../examples/lulesh/lulesh-init.cc         Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::Schedule<Kokkos::Static>, int>::R...         Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::Schedule<Kokkos::Static>, int>::R...
+           2506  ../examples/lulesh/lulesh-init.cc         Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::Schedule<Kokkos::Static>, long>::...         Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::Schedule<Kokkos::Static>, long>::...
+            271  ../examples/lulesh/lulesh-init.cc         Kokkos::StaticCrsGraph<int, Kokkos::LayoutLeft, Kokkos::OpenMP, Kokkos::Memor...         Kokkos::StaticCrsGraph<int, Kokkos::LayoutLeft, Kokkos::OpenMP, Kokkos::Memor...
+            470  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int* [8], Kokkos::LayoutRight>::View<std::__cxx11::basic_string<...         Kokkos::View<int* [8], Kokkos::LayoutRight>::View<std::__cxx11::basic_string<...
+            323  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int* [8], Kokkos::LayoutRight>::View<std::__cxx11::basic_string<...         Kokkos::View<int* [8], Kokkos::LayoutRight>::View<std::__cxx11::basic_string<...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*, Kokkos::HostSpace>::View<char [10]>                                   Kokkos::View<int*, Kokkos::HostSpace>::View<char [10]>(View<int *, Kokkos::Ho...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*, Kokkos::HostSpace>::View<char [14]>                                   Kokkos::View<int*, Kokkos::HostSpace>::View<char [14]>(View<int *, Kokkos::Ho...
+            462  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*, Kokkos::HostSpace>::View<std::__cxx11::basic_string<char, ...         Kokkos::View<int*, Kokkos::HostSpace>::View<std::__cxx11::basic_string<char, ...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<char [16]>                                                      Kokkos::View<int*>::View<char [16]>(View<int *> *, arg_label, type, const siz...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<char [19]>                                                      Kokkos::View<int*>::View<char [19]>(View<int *> *, arg_label, type, const siz...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<char [21]>                                                      Kokkos::View<int*>::View<char [21]>(View<int *> *, arg_label, type, const siz...
+            462  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<std::__cxx11::basic_string<char, std::char_traits<ch...         Kokkos::View<int*>::View<std::__cxx11::basic_string<char, std::char_traits<ch...
+            323  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<std::__cxx11::basic_string<char, std::char_traits<ch...         Kokkos::View<int*>::View<std::__cxx11::basic_string<char, std::char_traits<ch...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<double*, , double*, Kokkos::LayoutRight, Kokkos::Device<Kok...         Kokkos::deep_copy<double*, , double*, Kokkos::LayoutRight, Kokkos::Device<Kok...
+           1052  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<double*>                                                               Kokkos::deep_copy<double*>(dst, value) [lulesh-init.cc]
+           1050  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<double, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP,...         Kokkos::deep_copy<double, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP,...
+           7686  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenM...         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenM...
+           7686  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, int* [8], Kokkos::LayoutRigh...         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, int* [8], Kokkos::LayoutRigh...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int*, , int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::O...         Kokkos::deep_copy<int*, , int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::O...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int*, Kokkos::LayoutLeft, Kokkos::Device<Kokkos::OpenMP, Ko...         Kokkos::deep_copy<int*, Kokkos::LayoutLeft, Kokkos::Device<Kokkos::OpenMP, Ko...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP, K...         Kokkos::deep_copy<int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP, K...
+            863  ../examples/lulesh/lulesh-init.cc         Kokkos::impl_resize<, int* [8], Kokkos::LayoutRight>                                     type Kokkos::impl_resize<, int* [8], Kokkos::LayoutRight>(v, const size_t, co...
+            854  ../examples/lulesh/lulesh-init.cc         Kokkos::impl_resize<, int*>                                                              type Kokkos::impl_resize<, int*>(v, const size_t, const size_t, const size_t,...
+            697  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...
+            706  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...
+            912  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            791  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            791  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            944  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...
+            839  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...
+            126  ../examples/lulesh/lulesh-init.cc         _GLOBAL__sub_I_lulesh_init.cc                                                            _GLOBAL__sub_I_lulesh_init.cc() [lulesh-init.cc]
+           6589  ../examples/lulesh/lulesh-util.cc         Kokkos::deep_copy<double*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP...         Kokkos::deep_copy<double*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP...
+           1345  ../examples/lulesh/lulesh-util.cc         ParseCommandLineOptions                                                                  ParseCommandLineOptions(int, char * *, int, cmdLineOpts *) [lulesh-util.cc:67]
+            171  ../examples/lulesh/lulesh-util.cc         PrintCommandLineOptions                                                                  PrintCommandLineOptions(char *, int) [lulesh-util.cc:31]
+             67  ../examples/lulesh/lulesh-util.cc         StrToInt                                                                                 int StrToInt(const char *, int *) [lulesh-util.cc:13]
+            706  ../examples/lulesh/lulesh-util.cc         VerifyAndWriteFinalOutput                                                                VerifyAndWriteFinalOutput(Real_t, locDom, Int_t, Int_t) [lulesh-util.cc:222]
+            126  ../examples/lulesh/lulesh-util.cc         _GLOBAL__sub_I_lulesh_util.cc                                                            _GLOBAL__sub_I_lulesh_util.cc() [lulesh-util.cc]
+             17  ../examples/lulesh/lulesh-viz.cc          DumpToVisit                                                                              DumpToVisit(domain, int, int, int) [lulesh-viz.cc:415]
+            126  ../examples/lulesh/lulesh-viz.cc          _GLOBAL__sub_I_lulesh_viz.cc                                                             _GLOBAL__sub_I_lulesh_viz.cc() [lulesh-viz.cc]
+            451  ../examples/lulesh/lulesh.cc              .omp_outlined..103                                                                       .omp_outlined..103(const , const , const ParallelReduce<(lambda at ../example...
+            796  ../examples/lulesh/lulesh.cc              .omp_outlined..109                                                                       .omp_outlined..109(const , const , const ParallelFor<(lambda at ../examples/l...
+            394  ../examples/lulesh/lulesh.cc              .omp_outlined..111                                                                       .omp_outlined..111(const , const , const ParallelFor<(lambda at ../examples/l...
+            402  ../examples/lulesh/lulesh.cc              .omp_outlined..113                                                                       .omp_outlined..113(const , const , const ParallelFor<(lambda at ../examples/l...
+            427  ../examples/lulesh/lulesh.cc              .omp_outlined..115                                                                       .omp_outlined..115(const , const , const ParallelReduce<(lambda at ../example...
+            859  ../examples/lulesh/lulesh.cc              .omp_outlined..119                                                                       .omp_outlined..119(const , const , const ParallelFor<(lambda at ../examples/l...
+            243  ../examples/lulesh/lulesh.cc              .omp_outlined..122                                                                       .omp_outlined..122(const , const , const ParallelFor<(lambda at ../examples/l...
+            426  ../examples/lulesh/lulesh.cc              .omp_outlined..124                                                                       .omp_outlined..124(const , const , const ParallelFor<(lambda at ../examples/l...
+            529  ../examples/lulesh/lulesh.cc              .omp_outlined..127                                                                       .omp_outlined..127(const , const , const ParallelFor<(lambda at ../examples/l...
+            865  ../examples/lulesh/lulesh.cc              .omp_outlined..130                                                                       .omp_outlined..130(const , const , const ParallelFor<(lambda at ../examples/l...
+            539  ../examples/lulesh/lulesh.cc              .omp_outlined..132                                                                       .omp_outlined..132(const , const , const ParallelReduce<(lambda at ../example...
+            456  ../examples/lulesh/lulesh.cc              .omp_outlined..134                                                                       .omp_outlined..134(const , const , const ParallelReduce<(lambda at ../example...
+            252  ../examples/lulesh/lulesh.cc              .omp_outlined..20                                                                        .omp_outlined..20(const , const , const ParallelFor<(lambda at ../examples/lu...
+            870  ../examples/lulesh/lulesh.cc              .omp_outlined..35                                                                        .omp_outlined..35(const , const , const ParallelFor<(lambda at ../examples/lu...
+            473  ../examples/lulesh/lulesh.cc              .omp_outlined..42                                                                        .omp_outlined..42(const , const , const ParallelFor<(lambda at ../examples/lu...
+            252  ../examples/lulesh/lulesh.cc              .omp_outlined..46                                                                        .omp_outlined..46(const , const , const ParallelFor<(lambda at ../examples/lu...
+           1101  ../examples/lulesh/lulesh.cc              .omp_outlined..48                                                                        .omp_outlined..48(const , const , const ParallelFor<(lambda at ../examples/lu...
+            427  ../examples/lulesh/lulesh.cc              .omp_outlined..55                                                                        .omp_outlined..55(const , const , const ParallelReduce<(lambda at ../examples...
+           1326  ../examples/lulesh/lulesh.cc              .omp_outlined..57                                                                        .omp_outlined..57(const , const , const ParallelReduce<(lambda at ../examples...
+            243  ../examples/lulesh/lulesh.cc              .omp_outlined..61                                                                        .omp_outlined..61(const , const , const ParallelFor<(lambda at ../examples/lu...
+           1101  ../examples/lulesh/lulesh.cc              .omp_outlined..63                                                                        .omp_outlined..63(const , const , const ParallelFor<(lambda at ../examples/lu...
+            372  ../examples/lulesh/lulesh.cc              .omp_outlined..66                                                                        .omp_outlined..66(const , const , const ParallelFor<(lambda at ../examples/lu...
+            499  ../examples/lulesh/lulesh.cc              .omp_outlined..71                                                                        .omp_outlined..71(const , const , const ParallelFor<(lambda at ../examples/lu...
+            499  ../examples/lulesh/lulesh.cc              .omp_outlined..73                                                                        .omp_outlined..73(const , const , const ParallelFor<(lambda at ../examples/lu...
+            499  ../examples/lulesh/lulesh.cc              .omp_outlined..75                                                                        .omp_outlined..75(const , const , const ParallelFor<(lambda at ../examples/lu...
+            465  ../examples/lulesh/lulesh.cc              .omp_outlined..78                                                                        .omp_outlined..78(const , const , const ParallelFor<(lambda at ../examples/lu...
+            396  ../examples/lulesh/lulesh.cc              .omp_outlined..81                                                                        .omp_outlined..81(const , const , const ParallelFor<(lambda at ../examples/lu...
+            656  ../examples/lulesh/lulesh.cc              .omp_outlined..85                                                                        .omp_outlined..85(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            662  ../examples/lulesh/lulesh.cc              .omp_outlined..89                                                                        .omp_outlined..89(const , const , const ParallelFor<Kokkos::Impl::ViewCopy<Ko...
+            443  ../examples/lulesh/lulesh.cc              .omp_outlined..93                                                                        .omp_outlined..93(const , const , const ParallelReduce<(lambda at ../examples...
+            243  ../examples/lulesh/lulesh.cc              .omp_outlined..96                                                                        .omp_outlined..96(const , const , const ParallelFor<(lambda at ../examples/lu...
+            243  ../examples/lulesh/lulesh.cc              .omp_outlined..99                                                                        .omp_outlined..99(const , const , const ParallelFor<(lambda at ../examples/lu...
+          13367  ../examples/lulesh/lulesh.cc              ApplyMaterialPropertiesForElems                                                          ApplyMaterialPropertiesForElems(domain) [lulesh.cc:409]
+           1530  ../examples/lulesh/lulesh.cc              CalcElemCharacteristicLength                                                             Real_t CalcElemCharacteristicLength(const Real_t *, const Real_t *, const Rea...
+            982  ../examples/lulesh/lulesh.cc              CalcElemFBHourglassForce                                                                 CalcElemFBHourglassForce(const Real_t *, const Real_t[] *, coefficient, Real_...
+           2428  ../examples/lulesh/lulesh.cc              CalcElemNodeNormals                                                                      CalcElemNodeNormals(Real_t *, Real_t *, Real_t *, const Real_t *, const Real_...
+            853  ../examples/lulesh/lulesh.cc              CalcElemShapeFunctionDerivatives                                                         CalcElemShapeFunctionDerivatives(const Real_t *, const Real_t *, const Real_t...
+           1097  ../examples/lulesh/lulesh.cc              CalcElemVolumeDerivative                                                                 CalcElemVolumeDerivative(i, dvdx, dvdy, dvdz, const Real_t *, const Real_t *,...
+           1054  ../examples/lulesh/lulesh.cc              CalcKinematicsForElems                                                                   CalcKinematicsForElems(domain, Real_t, Index_t) [lulesh.cc]
+          14160  ../examples/lulesh/lulesh.cc              CalcVolumeForceForElems                                                                  CalcVolumeForceForElems(domain) [lulesh.cc:409]
+            366  ../examples/lulesh/lulesh.cc              Domain::AllocateGradients                                                                Domain::AllocateGradients(Domain *, Int_t, Int_t) [lulesh.cc:214]
+            475  ../examples/lulesh/lulesh.cc              Domain::DeallocateGradients                                                              Domain::DeallocateGradients(Domain *) [lulesh.cc:105]
+            250  ../examples/lulesh/lulesh.cc              Domain::DeallocateStrains                                                                Domain::DeallocateStrains(Domain *) [lulesh.cc:105]
+           4356  ../examples/lulesh/lulesh.cc              Domain::Domain                                                                           Domain::Domain(Domain *) [lulesh.cc:78]
+             15  ../examples/lulesh/lulesh.cc              Domain::delv_eta                                                                         Domain::delv_eta(const Domain *, const Index_t) [lulesh.cc:371]
+             15  ../examples/lulesh/lulesh.cc              Domain::delv_xi                                                                          Domain::delv_xi(const Domain *, const Index_t) [lulesh.cc:368]
+             15  ../examples/lulesh/lulesh.cc              Domain::delv_zeta                                                                        Domain::delv_zeta(const Domain *, const Index_t) [lulesh.cc:374]
+             15  ../examples/lulesh/lulesh.cc              Domain::fx                                                                               Domain::fx(const Domain *, const Index_t) [lulesh.cc:303]
+             15  ../examples/lulesh/lulesh.cc              Domain::fy                                                                               Domain::fy(const Domain *, const Index_t) [lulesh.cc:306]
+             15  ../examples/lulesh/lulesh.cc              Domain::fz                                                                               Domain::fz(const Domain *, const Index_t) [lulesh.cc:309]
+             15  ../examples/lulesh/lulesh.cc              Domain::nodalMass                                                                        Domain::nodalMass(const Domain *, const Index_t) [lulesh.cc:314]
+             15  ../examples/lulesh/lulesh.cc              Domain::x                                                                                Domain::x(const Domain *, const Index_t) [lulesh.cc:257]
+             15  ../examples/lulesh/lulesh.cc              Domain::xd                                                                               Domain::xd(const Domain *, const Index_t) [lulesh.cc:272]
+             15  ../examples/lulesh/lulesh.cc              Domain::y                                                                                Domain::y(const Domain *, const Index_t) [lulesh.cc:258]
+             15  ../examples/lulesh/lulesh.cc              Domain::yd                                                                               Domain::yd(const Domain *, const Index_t) [lulesh.cc:275]
+             15  ../examples/lulesh/lulesh.cc              Domain::z                                                                                Domain::z(const Domain *, const Index_t) [lulesh.cc:259]
+             15  ../examples/lulesh/lulesh.cc              Domain::zd                                                                               Domain::zd(const Domain *, const Index_t) [lulesh.cc:278]
+            330  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<doubl...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<doubl...
+            330  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<doubl...         Kokkos::Impl::ParallelConstructName<Kokkos::Impl::ViewCopy<Kokkos::View<doubl...
+           1508  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelFor<CalcEnergyForElems(double*, double*, double*, doubl...         type Kokkos::Impl::ParallelFor<CalcEnergyForElems(double*, double*, double*, ...
+           3606  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelFor<CalcFBHourglassForceForElems(Domain&, double*, Kokk...         type Kokkos::Impl::ParallelFor<CalcFBHourglassForceForElems(Domain&, double*,...
+           2917  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelFor<CalcKinematicsForElems(Domain&, double, int)::$_0, ...         type Kokkos::Impl::ParallelFor<CalcKinematicsForElems(Domain&, double, int)::...
+           3119  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelFor<CalcMonotonicQGradientsForElems(Domain&)::{lambda(i...         type Kokkos::Impl::ParallelFor<CalcMonotonicQGradientsForElems(Domain&)::{lam...
+           1969  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelFor<CalcMonotonicQRegionForElems(Domain&, int, double):...         type Kokkos::Impl::ParallelFor<CalcMonotonicQRegionForElems(Domain&, int, dou...
+           1265  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ParallelFor<IntegrateStressForElems(Domain&, double*, double*, ...         type Kokkos::Impl::ParallelFor<IntegrateStressForElems(Domain&, double*, doub...
+             49  ../examples/lulesh/lulesh.cc              Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewVal...         Kokkos::Impl::SharedAllocationRecord<Kokkos::HostSpace, Kokkos::Impl::ViewVal...
+           1497  ../examples/lulesh/lulesh.cc              Kokkos::Impl::TeamPolicyInternal<Kokkos::OpenMP>::TeamPolicyInternal                     Kokkos::Impl::TeamPolicyInternal<Kokkos::OpenMP>::TeamPolicyInternal(TeamPoli...
+            603  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewCopy<Kokkos::View<double*, Kokkos::LayoutLeft, Kokkos::Devi...         Kokkos::Impl::ViewCopy<Kokkos::View<double*, Kokkos::LayoutLeft, Kokkos::Devi...
+            604  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewCopy<Kokkos::View<double*, Kokkos::LayoutLeft, Kokkos::Devi...         Kokkos::Impl::ViewCopy<Kokkos::View<double*, Kokkos::LayoutLeft, Kokkos::Devi...
+            281  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            281  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...         Kokkos::Impl::ViewCtorProp<std::__cxx11::basic_string<char, std::char_traits<...
+            521  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewMapping<Kokkos::ViewTraits<double*>, void>::allocate_shared...         SharedAllocationRecord<void, void> * Kokkos::Impl::ViewMapping<Kokkos::ViewTr...
+            331  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewRemap<Kokkos::View<double*>, Kokkos::View<double*>, Kokkos:...         Kokkos::Impl::ViewRemap<Kokkos::View<double*>, Kokkos::View<double*>, Kokkos:...
+            461  ../examples/lulesh/lulesh.cc              Kokkos::Impl::ViewValueFunctor<Kokkos::Device<Kokkos::OpenMP, Kokkos::HostSpa...         enable_if_t<std::is_trivial<double>::value && std::is_trivially_copy_assignab...
+           1609  ../examples/lulesh/lulesh.cc              Kokkos::Impl::runtime_check_rank_host                                                    Kokkos::Impl::runtime_check_rank_host(const size_t, const bool, const size_t,...
+            697  ../examples/lulesh/lulesh.cc              Kokkos::Impl::view_copy<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::De...         Kokkos::Impl::view_copy<Kokkos::View<double*, Kokkos::LayoutRight, Kokkos::De...
+            697  ../examples/lulesh/lulesh.cc              Kokkos::Impl::view_copy<Kokkos::View<double*>, Kokkos::View<double*> >                   Kokkos::Impl::view_copy<Kokkos::View<double*>, Kokkos::View<double*> >(dst, s...
+           2250  ../examples/lulesh/lulesh.cc              Kokkos::RangePolicy<Kokkos::OpenMP>::RangePolicy                                         Kokkos::RangePolicy<Kokkos::OpenMP>::RangePolicy(RangePolicy<Kokkos::OpenMP> ...
+            213  ../examples/lulesh/lulesh.cc              Kokkos::StaticCrsGraph<int, Kokkos::LayoutLeft, Kokkos::OpenMP, Kokkos::Memor...         Kokkos::StaticCrsGraph<int, Kokkos::LayoutLeft, Kokkos::OpenMP, Kokkos::Memor...
+            410  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::View<char [6]>                                                    Kokkos::View<double*>::View<char [6]>(View<double *> *, arg_label, type, cons...
+            410  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::View<char [7]>                                                    Kokkos::View<double*>::View<char [7]>(View<double *> *, arg_label, type, cons...
+            462  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::View<std::__cxx11::basic_string<char, std::char_traits...         Kokkos::View<double*>::View<std::__cxx11::basic_string<char, std::char_traits...
+            323  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::View<std::__cxx11::basic_string<char, std::char_traits...         Kokkos::View<double*>::View<std::__cxx11::basic_string<char, std::char_traits...
+             25  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::~View                                                             Kokkos::View<double*>::~View(View<double *> *) [lulesh.cc:409]
+            840  ../examples/lulesh/lulesh.cc              Kokkos::abort                                                                            Kokkos::abort(const const char *, const const char *) [lulesh.cc:202]
+            854  ../examples/lulesh/lulesh.cc              Kokkos::impl_resize<, double*>                                                           type Kokkos::impl_resize<, double*>(v, const size_t, const size_t, const size...
+            928  ../examples/lulesh/lulesh.cc              Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            960  ../examples/lulesh/lulesh.cc              Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...
+          21470  ../examples/lulesh/lulesh.cc              LagrangeLeapFrog                                                                         LagrangeLeapFrog(domain) [lulesh.cc]
+            226  ../examples/lulesh/lulesh.cc              ResizeBuffer                                                                             ResizeBuffer(const size_t) [lulesh.cc:23]
+            169  ../examples/lulesh/lulesh.cc              _GLOBAL__sub_I_lulesh.cc                                                                 _GLOBAL__sub_I_lulesh.cc() [lulesh.cc]
+           1836  ../examples/lulesh/lulesh.cc              main                                                                                     int main(int, char * *) [lulesh.cc]
+             63  ../examples/lulesh/lulesh.cc              std::_Rb_tree<std::__cxx11::basic_string<char, std::char_traits<char>, std::a...         std::_Rb_tree<std::__cxx11::basic_string<char, std::char_traits<char>, std::a...
+             20  ../examples/lulesh/lulesh.cc              std::map<std::__cxx11::basic_string<char, std::char_traits<char>, std::alloca...         std::map<std::__cxx11::basic_string<char, std::char_traits<char>, std::alloca...
+            160  ../examples/lulesh/lulesh.cc              std::operator+<char, std::char_traits<char>, std::allocator<char> >                      basic_string<char, std::char_traits<char>, std::allocator<char> > std::operat...
+            187  ../examples/lulesh/lulesh.cc              std::pair<std::__cxx11::basic_string<char, std::char_traits<char>, std::alloc...         std::pair<std::__cxx11::basic_string<char, std::char_traits<char>, std::alloc...
+             11  lulesh                                    __clang_call_terminate                                                                   __clang_call_terminate() [lulesh]
+             33  lulesh                                    __do_global_dtors_aux                                                                    __do_global_dtors_aux() [lulesh]
+              5  lulesh                                    __libc_csu_fini                                                                          __libc_csu_fini() [lulesh]
+            101  lulesh                                    __libc_csu_init                                                                          __libc_csu_init() [lulesh]
+              5  lulesh                                    _dl_relocate_static_pie                                                                  _dl_relocate_static_pie() [lulesh]
+             13  lulesh                                    _fini                                                                                    _fini() [lulesh]
+             27  lulesh                                    _init                                                                                    _init() [lulesh]
+             47  lulesh                                    _start                                                                                   _start() [lulesh]
+              6  lulesh                                    frame_dummy                                                                              frame_dummy() [lulesh]
+
+An example of instrumented module and function info output
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+.. code-block:: shell
+
+   rocprof-sys-instrument -o lulesh.inst --label file line args --simulate -- lulesh
+
+After the heuristics are applied based on the pattern in :ref:`available-module-function-output`,
+the selected module and functions are:
+
+.. code-block:: shell
+
+   AddressRange  Module                                    Function                                                                                 FunctionSignature
+           9165  ../examples/lulesh/lulesh-comm.cc         CommMonoQ                                                                                CommMonoQ(domain) [lulesh-comm.cc:1891]
+           3396  ../examples/lulesh/lulesh-comm.cc         CommRecv                                                                                 CommRecv(domain, int, Index_t, Index_t, Index_t, Index_t, bool, bool) [lulesh...
+           8666  ../examples/lulesh/lulesh-comm.cc         CommSBN                                                                                  CommSBN(domain, int, Domain_member *) [lulesh-comm.cc:926]
+          10212  ../examples/lulesh/lulesh-comm.cc         CommSend                                                                                 CommSend(domain, int, Index_t, Domain_member *, Index_t, Index_t, Index_t, bo...
+           6823  ../examples/lulesh/lulesh-comm.cc         CommSyncPosVel                                                                           CommSyncPosVel(domain) [lulesh-comm.cc:1404]
+           1840  ../examples/lulesh/lulesh-init.cc         Domain::AllocateElemPersistent                                                           Domain::AllocateElemPersistent(Domain *, Int_t) [lulesh-init.cc:94]
+           1384  ../examples/lulesh/lulesh-init.cc         Domain::AllocateNodePersistent                                                           Domain::AllocateNodePersistent(Domain *, Int_t) [lulesh-init.cc:94]
+           1264  ../examples/lulesh/lulesh-init.cc         Domain::BuildMesh                                                                        Domain::BuildMesh(Domain *, Int_t, Int_t, Int_t) [lulesh-init.cc:308]
+           2312  ../examples/lulesh/lulesh-init.cc         Domain::CreateRegionIndexSets                                                            Domain::CreateRegionIndexSets(Domain *, Int_t, Int_t) [lulesh-init.cc:409]
+           7109  ../examples/lulesh/lulesh-init.cc         Domain::Domain                                                                           Domain::Domain(Domain *, Int_t, Index_t, Index_t, Index_t, Index_t, int, int,...
+           2458  ../examples/lulesh/lulesh-init.cc         Domain::SetupBoundaryConditions                                                          Domain::SetupBoundaryConditions(Domain *, Int_t) [lulesh-init.cc:409]
+            956  ../examples/lulesh/lulesh-init.cc         Domain::SetupCommBuffers                                                                 Domain::SetupCommBuffers(Domain *, Int_t) [lulesh-init.cc]
+           1456  ../examples/lulesh/lulesh-init.cc         Domain::SetupElementConnectivities                                                       Domain::SetupElementConnectivities(Domain *, Int_t) [lulesh-init.cc:409]
+            721  ../examples/lulesh/lulesh-init.cc         Domain::SetupSymmetryPlanes                                                              Domain::SetupSymmetryPlanes(Domain *, Int_t) [lulesh-init.cc:409]
+           1591  ../examples/lulesh/lulesh-init.cc         Domain::SetupThreadSupportStructures                                                     Domain::SetupThreadSupportStructures(Domain *) [lulesh-init.cc:376]
+           1644  ../examples/lulesh/lulesh-init.cc         Domain::~Domain                                                                          Domain::~Domain(Domain *) [lulesh-init.cc:286]
+            271  ../examples/lulesh/lulesh-init.cc         Kokkos::StaticCrsGraph<int, Kokkos::LayoutLeft, Kokkos::OpenMP, Kokkos::Memor...         Kokkos::StaticCrsGraph<int, Kokkos::LayoutLeft, Kokkos::OpenMP, Kokkos::Memor...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*, Kokkos::HostSpace>::View<char [10]>                                   Kokkos::View<int*, Kokkos::HostSpace>::View<char [10]>(View<int *, Kokkos::Ho...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*, Kokkos::HostSpace>::View<char [14]>                                   Kokkos::View<int*, Kokkos::HostSpace>::View<char [14]>(View<int *, Kokkos::Ho...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<char [16]>                                                      Kokkos::View<int*>::View<char [16]>(View<int *> *, arg_label, type, const siz...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<char [19]>                                                      Kokkos::View<int*>::View<char [19]>(View<int *> *, arg_label, type, const siz...
+            410  ../examples/lulesh/lulesh-init.cc         Kokkos::View<int*>::View<char [21]>                                                      Kokkos::View<int*>::View<char [21]>(View<int *> *, arg_label, type, const siz...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<double*, , double*, Kokkos::LayoutRight, Kokkos::Device<Kok...         Kokkos::deep_copy<double*, , double*, Kokkos::LayoutRight, Kokkos::Device<Kok...
+           1052  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<double*>                                                               Kokkos::deep_copy<double*>(dst, value) [lulesh-init.cc]
+           1050  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<double, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP,...         Kokkos::deep_copy<double, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP,...
+           7686  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenM...         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenM...
+           7686  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, int* [8], Kokkos::LayoutRigh...         Kokkos::deep_copy<int* [8], Kokkos::LayoutRight, int* [8], Kokkos::LayoutRigh...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int*, , int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::O...         Kokkos::deep_copy<int*, , int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::O...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int*, Kokkos::LayoutLeft, Kokkos::Device<Kokkos::OpenMP, Ko...         Kokkos::deep_copy<int*, Kokkos::LayoutLeft, Kokkos::Device<Kokkos::OpenMP, Ko...
+           6589  ../examples/lulesh/lulesh-init.cc         Kokkos::deep_copy<int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP, K...         Kokkos::deep_copy<int*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP, K...
+            697  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...
+            706  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...         Kokkos::parallel_for<Kokkos::MDRangePolicy<Kokkos::OpenMP, Kokkos::Rank<2u, (...
+            912  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            791  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            791  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            944  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...
+            839  ../examples/lulesh/lulesh-init.cc         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...
+           6589  ../examples/lulesh/lulesh-util.cc         Kokkos::deep_copy<double*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP...         Kokkos::deep_copy<double*, Kokkos::LayoutRight, Kokkos::Device<Kokkos::OpenMP...
+           1345  ../examples/lulesh/lulesh-util.cc         ParseCommandLineOptions                                                                  ParseCommandLineOptions(int, char * *, int, cmdLineOpts *) [lulesh-util.cc:67]
+            706  ../examples/lulesh/lulesh-util.cc         VerifyAndWriteFinalOutput                                                                VerifyAndWriteFinalOutput(Real_t, locDom, Int_t, Int_t) [lulesh-util.cc:222]
+          13367  ../examples/lulesh/lulesh.cc              ApplyMaterialPropertiesForElems                                                          ApplyMaterialPropertiesForElems(domain) [lulesh.cc:409]
+            982  ../examples/lulesh/lulesh.cc              CalcElemFBHourglassForce                                                                 CalcElemFBHourglassForce(const Real_t *, const Real_t[] *, coefficient, Real_...
+           2428  ../examples/lulesh/lulesh.cc              CalcElemNodeNormals                                                                      CalcElemNodeNormals(Real_t *, Real_t *, Real_t *, const Real_t *, const Real_...
+            853  ../examples/lulesh/lulesh.cc              CalcElemShapeFunctionDerivatives                                                         CalcElemShapeFunctionDerivatives(const Real_t *, const Real_t *, const Real_t...
+           1054  ../examples/lulesh/lulesh.cc              CalcKinematicsForElems                                                                   CalcKinematicsForElems(domain, Real_t, Index_t) [lulesh.cc]
+          14160  ../examples/lulesh/lulesh.cc              CalcVolumeForceForElems                                                                  CalcVolumeForceForElems(domain) [lulesh.cc:409]
+            366  ../examples/lulesh/lulesh.cc              Domain::AllocateGradients                                                                Domain::AllocateGradients(Domain *, Int_t, Int_t) [lulesh.cc:214]
+            475  ../examples/lulesh/lulesh.cc              Domain::DeallocateGradients                                                              Domain::DeallocateGradients(Domain *) [lulesh.cc:105]
+           4356  ../examples/lulesh/lulesh.cc              Domain::Domain                                                                           Domain::Domain(Domain *) [lulesh.cc:78]
+            410  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::View<char [6]>                                                    Kokkos::View<double*>::View<char [6]>(View<double *> *, arg_label, type, cons...
+            410  ../examples/lulesh/lulesh.cc              Kokkos::View<double*>::View<char [7]>                                                    Kokkos::View<double*>::View<char [7]>(View<double *> *, arg_label, type, cons...
+            928  ../examples/lulesh/lulesh.cc              Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<in...
+            960  ../examples/lulesh/lulesh.cc              Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...         Kokkos::parallel_for<Kokkos::RangePolicy<Kokkos::OpenMP, Kokkos::IndexType<lo...
+          21470  ../examples/lulesh/lulesh.cc              LagrangeLeapFrog                                                                         LagrangeLeapFrog(domain) [lulesh.cc]
+           1836  ../examples/lulesh/lulesh.cc              main                                                                                     int main(int, char * *) [lulesh.cc]
+
+Sampling
+========================================
+
+.. note::
+
+   This capability has been deprecated in favor of :doc:`Call stack sampling <./sampling-call-stack>`.
+
+By default, ``rocprof-sys-instrument`` uses ``--mode trace`` for instrumentation. The ``--mode sampling`` option
+only instruments ``main`` in an executable. It activates both CPU call-stack sampling and
+background system-level thread sampling by default.
+Tracing capabilities which do not rely on instrumentation, such as the HIP API and kernel tracing
+are still available.
+
+The ROCm Systems Profiler sampling capabilities are always available, even in trace mode, but are deactivated by default.
+To activate sampling in trace mode, set ``ROCPROFSYS_USE_SAMPLING=ON`` in the environment
+or in an ROCm Systems Profiler configuration file.
+
+Embedding a default configuration
+========================================
+
+Use the ``--env`` option to embed a default configuration into the target. Although this option
+works for runtime instrumentation, it is most useful when generating new binaries because the generated
+binary can be used later on in a different login session when the environment might have changed.
+
+For example, if the following commands are run,
+the configuration settings are not be preserved for subsequent sessions:
+
+.. code-block:: shell
+
+   rocprof-sys-instrument -o ./foo.inst -- ./foo
+   export ROCPROFSYS_USE_SAMPLING=ON
+   export ROCPROFSYS_SAMPLING_FREQ=5
+   rocprof-sys-run -- ./foo.inst
+
+Whereas the following command preserves those environment variables:
+
+.. code-block:: shell
+
+   rocprof-sys-instrument -o ./foo.samp --env ROCPROFSYS_USE_SAMPLING=ON ROCPROFSYS_SAMPLING_FREQ=5 -- ./foo
+
+They can now be used in future sessions.
+
+.. code-block:: shell
+
+   # will sample 5x per second
+   rocprof-sys-run -- ./foo.samp
+
+Even though the environment variables are preserved, subsequent sessions can still override those defaults:
+
+.. code-block:: shell
+
+   # will sample 100x per second
+   export ROCPROFSYS_SAMPLING_FREQ=100
+   rocprof-sys-run -- ./foo.samp
+
+.. _rpath-troubleshooting:
+
+Troubleshooting
+----------------------------------------------
+
+Checking for RPATH
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+If ``ldd ./foo.inst`` from the :ref:`binary-rewriting-library-label`
+section still returns ``/usr/local/lib/libfoo.so.2``, the executable could have
+an rpath encoded in the binary.
+This ELF entry results in the dynamic linker ignoring ``LD_LIBRARY_PATH`` if
+it finds ``libfoo.so.2`` in the rpath.
+Using the ``objdump`` tool, perform the following query:
+
+.. code-block:: shell
+
+   objdump -p <exe-or-library> | egrep 'RPATH|RUNPATH'
+
+If this produces output that appears similar to this output.:
+
+.. code-block:: shell
+
+   RUNPATH              $ORIGIN:$ORIGIN/../lib
+
+Remove or modify the rpath to get ``foo.inst`` to resolve
+to the instrumented ``libfoo.so.2`` as explained in the next section.
+
+Modifying an RPATH
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+This code snippet uses the ``patchelf`` tool to modify the rpath of the given executable
+or library to ``/home/user``, which is where the instrumented libraries are located.
+
+.. note::
+
+   This functionality requires the ``patchelf`` package.
+
+.. code-block:: shell
+
+   patchelf --remove-rpath <exe-or-library>
+   patchelf --set-rpath '/home/user' <exe-or-library>

projects/rocprofiler-systems/docs/how-to/nic-profiling.rst

@@ -0,0 +1,131 @@
+.. meta::
+   :description: ROCm Systems Profiler network performance profiling
+   :keywords: rocprof-sys, rocprofiler-systems, ROCm, tips, how to, profiler, tracking, NIC, network, AMD
+
+********************************************
+Network performance profiling
+********************************************
+
+`ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ supports network profiling.
+
+All network events that can be traced on the system can be listed by running the command:
+
+.. code-block:: shell
+
+    rocprof-sys-avail -H -r net
+
+For example, if the system's NIC is enp7s0, then the output of this command looks like:
+
+.. code-block:: shell
+
+  |-------------------------------|---------|-----------|-------------------------------|
+  |       HARDWARE COUNTER        | DEVICE  | AVAILABLE |            SUMMARY            |
+  |-------------------------------|---------|-----------|-------------------------------|
+  | net:::enp7s0:rx:byte          |   CPU   |   true    | enp7s0 receive byte           |
+  | net:::enp7s0:rx:packet        |   CPU   |   true    | enp7s0 receive packet         |
+  | net:::enp7s0:rx:error         |   CPU   |   true    | enp7s0 receive error          |
+  | net:::enp7s0:rx:droppe        |   CPU   |   true    | enp7s0 receive droppe         |
+  | net:::enp7s0:rx:fif           |   CPU   |   true    | enp7s0 receive fif            |
+  | net:::enp7s0:rx:fram          |   CPU   |   true    | enp7s0 receive fram           |
+  | net:::enp7s0:rx:compresse     |   CPU   |   true    | enp7s0 receive compresse      |
+  | net:::enp7s0:rx:multicas      |   CPU   |   true    | enp7s0 receive multicas       |
+  | net:::enp7s0:tx:byte          |   CPU   |   true    | enp7s0 transmit byte          |
+  | net:::enp7s0:tx:packet        |   CPU   |   true    | enp7s0 transmit packet        |
+  | net:::enp7s0:tx:error         |   CPU   |   true    | enp7s0 transmit error         |
+  | net:::enp7s0:tx:droppe        |   CPU   |   true    | enp7s0 transmit droppe        |
+  | net:::enp7s0:tx:fif           |   CPU   |   true    | enp7s0 transmit fif           |
+  | net:::enp7s0:tx:coll          |   CPU   |   true    | enp7s0 transmit coll          |
+  | net:::enp7s0:tx:carrie        |   CPU   |   true    | enp7s0 transmit carrie        |
+  | net:::enp7s0:tx:compresse     |   CPU   |   true    | enp7s0 transmit compresse     |
+  |-------------------------------|---------|-----------|-------------------------------|
+
+To track bytes and packets sent and received by the NIC ``enp7s0``, the configuration parameters should be configured as the following example:
+
+.. code-block:: shell
+
+  ROCPROFSYS_PAPI_EVENTS = net:::enp7s0:tx:byte net:::enp7s0:rx:byte net:::enp7s0:tx:packet net:::enp7s0:rx:packet
+
+Configuration
+=============
+
+A sample configuration parameter settings looks like:
+
+.. code-block:: shell
+
+  ROCPROFSYS_SAMPLING_FREQ=10
+  ROCPROFSYS_USE_SAMPLING=ON
+  ROCPROFSYS_TIMEMORY_COMPONENTS=wall_clock papi_array network_stats
+  ROCPROFSYS_NETWORK_INTERFACE=enp7s0
+  ROCPROFSYS_PAPI_EVENTS=net:::enp7s0:tx:byte net:::enp7s0:rx:byte net:::enp7s0:rx:packet net:::enp7s0:tx:packet
+
+Details of the configuration parameter settings configured in the example are:
+
+* **Sampling Frequency**: 10 samples per second
+* **TIMEMORY**:  Outputs the summaries for the ``wall_clock``, ``papi_array``, and ``network_stats`` components.
+* **Network Interface**: ``enp7s0`` is the predictable network interface device name.
+* **Events for the network device to be sampled**: Bytes transmitted, bytes received, packets transmitted, and packets received.
+
+The configuration parameter settings can be saved in a configuration file. Here is an example of a complete configuration file, ``rocprofsys.cfg``:
+
+.. code-block:: shell
+
+  ROCPROFSYS_VERBOSE=1
+  ROCPROFSYS_DL_VERBOSE=1
+  ROCPROFSYS_SAMPLING_FREQ=10
+  ROCPROFSYS_SAMPLING_DELAY=0.05
+  ROCPROFSYS_SAMPLING_CPUS=0-9
+  ROCPROFSYS_SAMPLING_GPUS=$env:HIP_VISIBLE_DEVICES
+  ROCPROFSYS_TRACE=ON
+  ROCPROFSYS_PROFILE=ON
+  ROCPROFSYS_USE_SAMPLING=ON
+  ROCPROFSYS_USE_PROCESS_SAMPLING=OFF
+  ROCPROFSYS_TIME_OUTPUT=OFF
+  ROCPROFSYS_FILE_OUTPUT=ON
+  ROCPROFSYS_TIMEMORY_COMPONENTS=wall_clock papi_array network_stats
+  ROCPROFSYS_USE_PID=OFF
+  ROCPROFSYS_OUTPUT_PREFIX=foo/
+  ROCPROFSYS_NETWORK_INTERFACE=enp7s0
+  ROCPROFSYS_PAPI_EVENTS = net:::enp7s0:tx:byte net:::enp7s0:rx:byte net:::enp7s0:rx:packet net:::enp7s0:tx:packet
+
+To specify the configuration file, use the ``ROCPROFSYS_CONFIG_FILE`` setting:
+
+.. code-block:: shell
+
+  ROCPROFSYS_CONFIG_FILE=/path/to/rocprofsys.cfg
+
+This setting defines the location of the ROCm Systems Profiler configuration file.
+
+.. note::
+
+   To collect network counters using Process Application Program Interface (PAPI), ensure that
+   `/proc/sys/kernel/perf_event_paranoid` has a value <= 2. See
+   :ref:`rocprof-sys_papi_events`
+   for details.
+
+Instrumenting and running a program
+===================================
+
+An example rocprof-sys-instrument command is:
+
+.. code-block:: shell
+
+  rocprof-sys-instrument -o foo.inst  \
+    --log-file mylog.log --verbose --debug \
+    "--print-instrumented" "functions" "-e" "-v" "2" "--caller-include" \
+    "inner" "-i" "4096" "--" ./foo
+
+This command generates an instrumented binary ``foo.inst``. Then, run
+it with the following command:
+
+.. code-block:: shell
+
+  rocprof-sys-run -- ./foo.inst
+
+To view the generated ``.proto`` file in the browser, open the
+`Perfetto UI page <https://ui.perfetto.dev/>`_. Then, click on
+``Open trace file`` and select the ``.proto`` file. In the browser, it looks
+like this:
+
+.. image:: ../data/rocprof-sys-perfetto-nic-trace.png
+   :alt: Visualization of a performance graph in Perfetto with network tracks
+   :width: 800

projects/rocprofiler-systems/docs/how-to/performing-causal-profiling.rst

@@ -0,0 +1,626 @@
+.. meta::
+   :description: ROCm Systems Profiler causal profiling documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, causal profiling, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Performing causal profiling
+****************************************************
+
+The process of causal profiling can be summarized as:
+
+*If you speed up a given block of code by X%, the application will run Y% faster*.
+
+Causal profiling directs parallel application developers to where they should focus their optimization
+efforts by quantifying the potential impact of optimizations. Causal profiling is rooted in the concept
+that *software execution speed is relative*. Speeding up a block of code by X% is mathematically equivalent
+to that block of code running at its current speed if all the other code is running slower by X%.
+Thus, causal profiling works by performing experiments on blocks of code during program execution which
+insert pauses to slow down all other concurrently running code. During post-processing, these experiments
+are translated into calculations for the potential impact of speeding up this block of code.
+
+Consider the following C++ code executing ``foo`` and ``bar`` concurrently in two different threads
+where ``foo`` is ideally 30% faster than ``bar``:
+
+.. code-block:: cpp
+
+   #include <cstddef>
+   #include <thread>
+   constexpr size_t FOO_N =  7 * 1000000000UL;
+   constexpr size_t BAR_N = 10 * 1000000000UL;
+
+   void foo()
+   {
+      for(volatile size_t i = 0; i < FOO_N; ++i) {}
+   }
+
+   void bar()
+   {
+      for(volatile size_t i = 0; i < BAR_N; ++i) {}
+   }
+
+   int main()
+   {
+      std::thread _threads[] = { std::thread{ foo },
+                        std::thread{ bar } };
+
+      for(auto& itr : _threads)
+         itr.join();
+   }
+
+No matter how many optimizations are applied to ``foo``, the application will always
+require the same amount of time
+because the end-to-end performance is limited by ``bar``. However, a 5% speed-up
+in ``bar`` results in the
+end-to-end performance improving by 5%. This trend continues linearly, with a 10% speed-up
+in ``bar`` yielding a 10% speed-up in
+end-to-end performance, and so on, up to a 30% speed-up, at which point ``bar`` runs as fast as ``foo``.
+Any speed-up to ``bar`` beyond 30% still only yields an end-to-end performance
+improvement of 30% because the application
+is now limited by performance of ``foo``, as demonstrated below in the causal
+profiling visualization:
+
+.. image:: ../data/causal-foobar.png
+   :alt: Visualization of the performance improvements for two functions with causal profiling
+
+The full details of the causal profiling methodology can be found in the paper
+`Coz: Finding Code that Counts with Causal Profiling <http://arxiv.org/pdf/1608.03676v1.pdf>`_.
+The author's implementation is publicly available on `GitHub <https://github.com/plasma-umass/coz>`_.
+
+Getting started
+========================================
+
+To effectively use causal profiling, it is important to understand a few key
+concepts, such as progress points.
+
+Progress points
+-----------------------------------
+
+Causal profiling requires "progress points" to track progress through the code
+in between samples. Progress points must be triggered in a deterministic manner via instrumentation.
+This can happen in three different ways:
+
+* `ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ can leverage the callbacks from
+  Kokkos-Tools, OpenMP-Tools, rocprofiler-sdk, etc. and the wrappers around functions for
+  MPI, NUMA, RCCL, etc. to act as progress points
+* Users can leverage the :doc:`runtime instrumentation capabilities <./instrumenting-rewriting-binary-application>`
+  to insert progress points
+* Users can leverage :doc:`User APIs <../how-to/using-rocprof-sys-api>`,
+  such as ``ROCPROFSYS_CAUSAL_PROGRESS``
+
+.. note::
+
+   Binary rewrite to insert progress points is not supported. When a rewritten binary
+   runs, Dyninst translates the instruction pointer address in order to perform
+   the instrumentation. As a result, call stack samples never return instruction
+   pointer addresses within the valid ROCm Systems Profiler range.
+
+Key concepts
+-----------------------------------
+
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| Concept          | Setting                              | Options                          | Description                                |
++==================+======================================+==================================+============================================+
+| Backend          | ``ROCPROFSYS_CAUSAL_BACKEND``        | ``perf``, ``timer``              | Backend for recording samples required     |
+|                  |                                      |                                  | to calculate the virtual speed-up          |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| Mode             | ``ROCPROFSYS_CAUSAL_MODE``           | ``function``, ``line``           | Select an entire function or individual    |
+|                  |                                      |                                  | line of code for causal experiments        |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| End-to-end       | ``ROCPROFSYS_CAUSAL_END_TO_END``     | Boolean                          | Perform a single experiment during the     |
+|                  |                                      |                                  | entire run (does not require               |
+|                  |                                      |                                  | progress points)                           |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| Fixed speed-up   | ``ROCPROFSYS_CAUSAL_FIXED_SPEEDUP``  | one or more values from [0, 100] | Virtual speed-up or pool of virtual        |
+|                  |                                      |                                  | speed-ups to randomly select               |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| Binary scope     | ``ROCPROFSYS_CAUSAL_BINARY_SCOPE``   | regular expression(s)            | Dynamic binaries containing code for       |
+|                  |                                      |                                  | experiments                                |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| Source scope     | ``ROCPROFSYS_CAUSAL_SOURCE_SCOPE``   | regular expression(s)            | ``<file>`` and/or ``<file>:<line>``        |
+|                  |                                      |                                  | containing code to include in experiments  |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+| Function scope   | ``ROCPROFSYS_CAUSAL_FUNCTION_SCOPE`` | regular expression(s)            | Restricts experiments to matching          |
+|                  |                                      |                                  | functions (function mode) or lines of      |
+|                  |                                      |                                  | code within matching functions (line mode) |
++------------------+--------------------------------------+----------------------------------+--------------------------------------------+
+
+.. note::
+
+   * Binary scope defaults to ``%MAIN%`` (in the executable), but the scope can be expanded to include linked libraries.
+   * ``<file>`` and ``<file>:<line>`` support requires debug info (for example, the code must be compiled with ``-g`` or, preferably, with ``-g3``)
+   * Function mode does not require debug info but does not support stripped binaries
+
+Backends
+-----------------------------------
+
+There are two backends to choose from: ``perf`` and ``timer``.
+They are used to record the samples required to calculate the virtual speedup.
+Both backends interrupt each thread 1000 times per second (of CPU-time) to apply the virtual speed-ups.
+The difference between each backend is how the samples are recorded.
+There are three key differences between the two backends:
+
+* the ``perf`` backend requires Linux Perf and elevated security priviledges
+* the ``perf`` backend interrupts the application less frequently whereas the ``timer`` backend
+  interrupts the application 1000 times per second of realtime
+* the ``timer`` backend has less accurate call stacks due to instruction pointer skid
+
+In general, the ``perf`` backend is preferred over the ``timer`` backend when sufficient
+security priviledges permit its usage.
+If ``ROCPROFSYS_CAUSAL_BACKEND`` is set to ``auto``, ROCm Systems Profiler falls back
+to using the ``timer`` backend only if
+the ``perf`` backend fails. If ``ROCPROFSYS_CAUSAL_BACKEND`` is
+set to ``perf`` and using this backend fails, ROCm Systems Profiler aborts.
+
+Instruction pointer skid
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Instruction pointer (IP) skid measures how many instructions run after the event of interest
+before the program actually stops. The IP skid is calculated by subtracting
+the location of the IP at the point of interest from the location of the IP
+when the kernel finally stops the application.
+For the ``timer`` backend, this translates to the
+difference in the IP between when the timer generated a signal and when the
+signal was actually generated. Although IP skid still occurs with the ``perf`` backend,
+it is much more pronounced with the ``timer`` backend due to the overhead of pausing the entire thread.
+This means the ``timer`` backend tends to have a lower resolution than the ``perf`` backend,
+especially in ``line`` mode.
+
+Installing Linux Perf
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Linux Perf is built into the kernel and may already be installed
+(for instance, it is included in the default kernel for OpenSUSE).
+The official method of checking whether Linux Perf is installed is
+checking for the existence of the file
+``/proc/sys/kernel/perf_event_paranoid``. If the file exists, the kernel has Perf installed.
+
+If this file does not exist, as with Debian-based systems like Ubuntu, run the following command as superuser:
+
+.. code-block:: shell
+
+   apt-get install linux-tools-common linux-tools-generic linux-tools-$(uname -r)
+
+and reboot your computer. In order to use the ``perf`` backend, the value
+of ``/proc/sys/kernel/perf_event_paranoid``
+should be less than or equal to 2. If the value in this file is greater than 2, you can't
+use the ``perf`` backend.
+
+To update the paranoid level temporarily until the system is rebooted, run
+one of the following commands
+as a superuser (where ``PARANOID_LEVEL=<N>`` has a value of ``<N>`` in the range ``[-1, 2]``):
+
+.. code-block:: shell
+
+   echo ${PARANOID_LEVEL} | sudo tee /proc/sys/kernel/perf_event_paranoid
+
+or
+
+.. code-block:: shell
+
+   sysctl kernel.perf_event_paranoid=${PARANOID_LEVEL}
+
+To make the paranoid level persistent after a reboot, add ``kernel.perf_event_paranoid=<N>``
+(where ``<N>`` is the desired paranoid level) to the ``/etc/sysctl.conf`` file.
+
+Speed-up prediction variability and the rocprof-sys-causal executable
+-----------------------------------------------------------------------
+
+Causal profiling typically requires running the application several times in
+order to adequately sample all the code domains, experiment
+with speed-ups and other techniques, and resolve statistical fluctuations.
+The ``rocprof-sys-causal`` executable is designed to simplify this procedure:
+
+.. code-block:: shell
+
+   $ rocprof-sys-causal --help
+   [rocprof-sys-causal] Usage: ./bin/rocprof-sys-causal [ --help (count: 0, dtype: bool)
+                                                      --version (count: 0, dtype: bool)
+                                                      --monochrome (max: 1, dtype: bool)
+                                                      --debug (max: 1, dtype: bool)
+                                                      --verbose (count: 1)
+                                                      --config (min: 0, dtype: filepath)
+                                                      --launcher (count: 1, dtype: executable)
+                                                      --generate-configs (min: 0, dtype: folder)
+                                                      --no-defaults (min: 0, dtype: bool)
+                                                      --mode (count: 1, dtype: string)
+                                                      --output-name (min: 1, dtype: filename)
+                                                      --reset (max: 1, dtype: bool)
+                                                      --end-to-end (max: 1, dtype: bool)
+                                                      --wait (count: 1, dtype: seconds)
+                                                      --duration (count: 1, dtype: seconds)
+                                                      --iterations (count: 1, dtype: int)
+                                                      --speedups (min: 0, dtype: integers)
+                                                      --binary-scope (min: 0, dtype: integers)
+                                                      --source-scope (min: 0, dtype: integers)
+                                                      --function-scope (min: 0, dtype: regex-list)
+                                                      --binary-exclude (min: 0, dtype: integers)
+                                                      --source-exclude (min: 0, dtype: integers)
+                                                      --function-exclude (min: 0, dtype: regex-list)
+                                                   ]
+
+      Causal profiling usually requires multiple runs to reliably resolve the speedup estimates.
+      This executable is designed to streamline that process.
+      For example (assume all commands end with \'-- <exe> <args>\'):
+
+         rocprof-sys-causal -n 5 -- <exe>                  # runs <exe> 5x with causal profiling enabled
+
+         rocprof-sys-causal -s 0 5,10,15,20                # runs <exe> 2x with virtual speedups:
+                                                         #   - 0
+                                                         #   - randomly selected from 5, 10, 15, and 20
+
+         rocprof-sys-causal -F func_A func_B func_(A|B)    # runs <exe> 3x with the function scope limited to:
+                                                         #   1. func_A
+                                                         #   2. func_B
+                                                         #   3. func_A or func_B
+      General tips:
+      - Insert progress points at hotspots in your code or use rocprof-sys\'s runtime instrumentation
+         - Note: binary rewrite will produce a incompatible new binary
+      - Run rocprof-sys-causal in "function" mode first (does not require debug info)
+      - Run rocprof-sys-causal in "line" mode when you are targeting one function (requires debug info)
+         - Preferably, use predictions from the "function" mode to determine which function to target
+      - Limit the virtual speedups to a smaller pool, e.g., 0,5,10,25,50, to get reliable predictions quicker
+      - Make use of the binary, source, and function scope to limit the functions/lines selected for experiments
+         - Note: source scope requires debug info
+
+
+   Options:
+      -h, -?, --help                 Shows this page
+      --version                      Prints the version and exit
+
+      [DEBUG OPTIONS]
+
+      --monochrome                   Disable colorized output
+      --debug                        Debug output
+      -v, --verbose                  Verbose output
+
+      [GENERAL OPTIONS]
+
+      -c, --config                   Base configuration file
+      -l, --launcher                 When running MPI jobs, rocprof-sys-causal needs to be *before* the executable which launches the MPI processes (i.e.
+                                    before `mpirun`, `srun`, etc.). Pass the name of the target executable (or a regex for matching to the name of the
+                                    target) for causal profiling, e.g., `rocprof-sys-causal -l foo -- mpirun -n 4 foo`. This ensures that the rocprof-sys
+                                    library is LD_PRELOADed on the proper target
+      -g, --generate-configs         Generate config files instead of passing environment variables directly. If no arguments are provided, the config files
+                                    will be placed in ${PWD}/rocprof-sys-causal-config folder
+      --no-defaults                  Do not activate default features which are recommended for causal profiling. For example: PID-tagging of output files
+                                    and timestamped subdirectories are disabled by default. Kokkos tools support is added by default
+                                    (ROCPROFSYS_USE_KOKKOSP=ON) because, for Kokkos applications, the Kokkos-Tools callbacks are used for progress points.
+                                    Activation of OpenMP tools support is similar
+
+      [CAUSAL PROFILING OPTIONS (General)]
+                                    (These settings will be applied to all causal profiling runs)
+
+      -m, --mode [ function (func) | line ]
+                                    Causal profiling mode
+      -o, --output-name              Output filename of causal profiling data w/o extension
+      -r, --reset                    Overwrite any existing experiment results during the first run
+      -e, --end-to-end               Single causal experiment for the entire application runtime
+      -w, --wait                     Set the wait time (i.e. delay) before starting the first causal experiment (in seconds)
+      -d, --duration                 Set the length of time (in seconds) to perform causal experimentationafter the first experiment is started. Once this
+                                    amount of time has elapsed, no more causal experiments will be started but any currently running experiment will be
+                                    allowed to finish.
+      -n, --iterations               Number of times to repeat the combination of run configurations
+
+      [CAUSAL PROFILING OPTIONS (Combinatorial)]
+                                    (Each individual argument to these options will multiply the number runs by the number of arguments and the number of
+                                    iterations. E.g. -n 2 -B "MAIN" -F "foo" "bar" will produce 4 runs: 2 iterations x 1 binary scope x 2 function scopes
+                                    (MAIN+foo, MAIN+bar, MAIN+foo, MAIN+bar))
+
+      -s, --speedups                 Pool of virtual speedups to sample from during experimentation. Each space designates a group and multiple speedups can
+                                    be grouped together by commas, e.g. -s 0 0,10,20-50 is two groups: group #1 is \'0\' and group #2 is \'0 10 20 25 30 35 40
+                                    45 50\'
+      -B, --binary-scope             Restricts causal experiments to the binaries matching the list of regular expressions. Each space designates a group
+                                    and multiple scopes can be grouped together with a semi-colon
+      -S, --source-scope             Restricts causal experiments to the source files or source file + lineno pairs (i.e. <file> or <file>:<line>) matching
+                                    the list of regular expressions. Each space designates a group and multiple scopes can be grouped together with a
+                                    semi-colon
+      -F, --function-scope           Restricts causal experiments to the functions matching the list of regular expressions. Each space designates a group
+                                    and multiple scopes can be grouped together with a semi-colon
+      -BE, --binary-exclude          Excludes causal experiments from being performed on the binaries matching the list of regular expressions. Each space
+                                    designates a group and multiple excludes can be grouped together with a semi-colon
+      -SE, --source-exclude          Excludes causal experiments from being performed on the code from the source files or source file + lineno pair (i.e.
+                                    <file> or <file>:<line>) matching the list of regular expressions. Each space designates a group and multiple excludes
+                                    can be grouped together with a semi-colon
+      -FE, --function-exclude        Excludes causal experiments from being performed on the functions matching the list of regular expressions. Each space
+                                    designates a group and multiple excludes can be grouped together with a semi-colon
+
+Examples
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+.. code-block:: shell
+
+   #!/bin/bash -e
+
+   module load rocprofiler-systems
+
+   N=20
+   I=3
+
+   # when providing speedups to rocprof-sys-causal, speedup
+   # groups are separated by a space so "0,10" results in
+   # one speedup group where rocprof-sys samples from
+   # the speedup set of {0, 10}. Passing "0 10" (without
+   # quotes to rocprof-sys-causal multiplies the
+   # number of runs by 2, where the first half of the
+   # runs instruct rocprof-sys to only use 0 as the
+   # speedup and the second half of the runs instruct
+   # rocprof-sys to only use 10 as the speedup.
+   SPEEDUPS="0,0,0,10,20,30,40,50,50,75,75,75,90,90,90"
+   # thus, -s ${SPEEDUPS} only multiplies the number
+   # of runs by 1 whereas -S ${SPEEDUPS_E2E} multiplies
+   # the number of runs by 15:
+   #   - 3 runs with speedup of 0
+   #   - 1 run for each of the speedups 10, 20, 30, and 40
+   #   - 2 runs with speedup of 50
+   #   - 3 runs with speedup of 75
+   #   - 3 runs with speedup of 90
+   SPEEDUPS_E2E=$(echo "${SPEEDUPS}" | sed \'s/,/ /g\')
+
+
+   # 20 iterations in function mode with 1 speedup group
+   # and source scope set to .cpp files
+   #
+   # outputs to files:
+   #   - causal/experiments.func.coz
+   #   - causal/experiments.func.json
+   #
+   # total executions: 20
+   #
+   rocprof-sys-causal     \
+      -n ${N}             \
+      -s ${SPEEDUPS}      \
+      -m function         \
+      -o experiments.func \
+      -S ".*\\.cpp"       \
+      --                  \
+      ./causal-rocprofsys-cpu "${@}"
+
+
+   # 20 iterations in line mode with 1 speedup group
+   # and source scope restricted to lines 100 and 110
+   # in the causal.cpp file.
+   #
+   # outputs to files:
+   #   - causal/experiments.line.coz
+   #   - causal/experiments.line.json
+   #
+   # total executions: 20
+   #
+   rocprof-sys-causal             \
+      -n ${N}                     \
+      -s ${SPEEDUPS}              \
+      -m line                     \
+      -o experiments.line         \
+      -S "causal\\.cpp:(100|110)" \
+      --                          \
+      ./causal-rocprofsys-cpu "${@}"
+
+
+   # 3 iterations in function mode of 15 singular speedups
+   # in end-to-end mode with 2 different function scopes
+   # where one is restricted to "cpu_slow_func" and
+   # another is restricted to "cpu_fast_func".
+   #
+   # outputs to files:
+   #   - causal/experiments.func.e2e.coz
+   #   - causal/experiments.func.e2e.json
+   #
+   # total executions: 90
+   #
+   rocprof-sys-causal         \
+      -n ${I}                 \
+      -s ${SPEEDUPS_E2E}      \
+      -m func                 \
+      -e                      \
+      -o experiments.func.e2e \
+      -F "cpu_slow_func"      \
+         "cpu_fast_func"      \
+      --                      \
+      ./causal-rocprofsys-cpu "${@}"
+
+   # 3 iterations in line mode of 15 singular speedups
+   # in end-to-end mode with 2 different source scopes
+   # where one is restricted to line 100 in causal.cpp
+   # and another is restricted to line 110 in causal.cpp.
+   #
+   # outputs to files:
+   #   - causal/experiments.line.e2e.coz
+   #   - causal/experiments.line.e2e.json
+   #
+   # total executions: 90
+   #
+   rocprof-sys-causal         \
+      -n ${I}                 \
+      -s ${SPEEDUPS_E2E}      \
+      -m line                 \
+      -e                      \
+      -o experiments.line.e2e \
+      -S "causal\\.cpp:100"   \
+         "causal\\.cpp:110"   \
+      --                      \
+      ./causal-rocprofsys-cpu "${@}"
+
+
+   export OMP_NUM_THREADS=8
+   export OMP_PROC_BIND=spread
+   export OMP_PLACES=threads
+
+   # set number of iterations to 5
+   N=5
+
+   # 5 iterations in function mode of 1 speedup
+   # group with the source scope restricted
+   # to files containing "lulesh" in their filename
+   # and exclude functions which start with "Kokkos::"
+   # or "std::enable_if".
+   #
+   # outputs to files:
+   #   - causal/experiments.func.coz
+   #   - causal/experiments.func.json
+   #
+   # total executions: 5
+   #
+   # First of 5 executions overwrites any
+   # existing causal/experiments.func.(coz|json)
+   # file due to "--reset" argument
+   #
+   rocprof-sys-causal                         \
+      --reset                                 \
+      -n ${N}                                 \
+      -s ${SPEEDUPS}                          \
+      -m func                                 \
+      -o experiments.func                     \
+      -S "lulesh.*"                           \
+      -FE "^(Kokkos::|std::enable_if)"        \
+      --                                      \
+      ./lulesh-rocprofsys -i 50 -s 200 -r 20 -b 5 -c 5 -p
+
+
+   # 5 iterations in line mode of 1 speedup
+   # group with the source scope restricted
+   # to files containing "lulesh" in their filename
+   # and exclude functions which start with "exec_range"
+   # or "execute" and which contain either
+   # "construct_shared_allocation" or "._omp_fn." in
+   # the function name.
+   #
+   # outputs to files:
+   #   - causal/experiments.line.coz
+   #   - causal/experiments.line.json
+   #
+   # total executions: 5
+   #
+   # First of 5 executions overwrites any
+   # existing causal/experiments.line.(coz|json)
+   # file due to "--reset" argument
+   #
+   rocprof-sys-causal                         \
+      --reset                                 \
+      -n ${N}                                 \
+      -s ${SPEEDUPS}                          \
+      -m line                                 \
+      -o experiments.line                     \
+      -S "lulesh.*"                           \
+      -FE "^(exec_range|execute);construct_shared_allocation;\\._omp_fn\\." \
+      --                                      \
+      ./lulesh-rocprofsys -i 50 -s 200 -r 20 -b 5 -c 5 -p
+
+
+   # 5 iterations in line mode of 1 speedup
+   # group with the source scope restricted
+   # to files whose basename is "lulesh.cc"
+   # for 3 different functions:
+   #   - ApplyMaterialPropertiesForElems
+   #   - CalcHourglassControlForElems
+   #   - CalcVolumeForceForElems
+   #
+   # outputs to files:
+   #   - causal/experiments.line.targeted.coz
+   #   - causal/experiments.line.targeted.json
+   #
+   # total executions: 15
+   #
+   # First of 5 executions overwrites any
+   # existing causal/experiments.line.(coz|json)
+   # file due to "--reset" argument
+   #
+   rocprof-sys-causal                            \
+      --reset                                 \
+      -n ${N}                                 \
+      -s ${SPEEDUPS}                          \
+      -m line                                 \
+      -o experiments.line.targeted            \
+      -F "ApplyMaterialPropertiesForElems"    \
+         "CalcHourglassControlForElems"       \
+         "CalcVolumeForceForElems"            \
+      -S "lulesh\\.cc"                        \
+      --                                      \
+      ./lulesh-rocprofsys -i 50 -s 200 -r 20 -b 5 -c 5 -p
+
+Using rocprof-sys-causal with other launchers like mpirun
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+The ``rocprof-sys-causal`` executable is intended to assist with application replay
+and is designed to always be at the start of the command line as the primary process.
+``rocprof-sys-causal`` typically adds a ``LD_PRELOAD`` of the ROCm Systems Profiler libraries
+into the environment before launching the command to inject the functionality
+required to start the causal profiling tooling. However, this is problematic
+when the target application for causal profiling uses a launcher, in which case
+it is listed as an argument rather than as the main application. For example,
+``foo`` is the target application for profiling, but the command to run it is
+``mpirun -n 2 foo``. Running the command ``rocprof-sys-causal -- mpirun -n 2 foo``
+applies the causal profiling to ``mpirun`` instead of ``foo``.
+
+``rocprof-sys-causal`` remedies this by providing a command-line option ``-l` / `--launcher``
+to indicate the target application is using a launcher script/executable. The
+argument to the command-line option is the name of, or regular expression for, the target application
+on the command line. When ``--launcher`` is used, ``rocprof-sys-causal`` generates
+all the replay configurations and runs them but delays adding the ``LD_PRELOAD``. Instead it
+inserts a call to itself into the command line right before the target
+application. This recursive call inherits the configuration from
+the parent ``rocprof-sys-causal`` executable, inserts an ``LD_PRELOAD`` into the environment,
+and calls ``execv`` to replace itself with the new process launched by the target
+application.
+
+In other words, the following command:
+
+.. code-block:: shell
+
+   rocprof-sys-causal -l foo -n 3 -- mpirun -n 2 foo`
+
+Effectively results in:
+
+.. code-block:: shell
+
+   mpirun -n 2 rocprof-sys-causal -- foo
+   mpirun -n 2 rocprof-sys-causal -- foo
+   mpirun -n 2 rocprof-sys-causal -- foo
+
+Visualizing the causal output
+-------------------------------------------------------------------------
+
+ROCm Systems Profiler generates ``causal/experiments.json`` and ``causal/experiments.coz`` in
+``${ROCPROFSYS_OUTPUT_PATH}/${ROCPROFSYS_OUTPUT_PREFIX}``. Visit
+`plasma-umass.org/coz <https://plasma-umass.org/coz/>`_ to open the ``*.coz`` file.
+
+ROCm Systems Profiler versus Coz
+=======================================
+
+This comparison is intended for readers who are familiar with the
+`Coz profiler <https://github.com/plasma-umass/coz>`_.
+ROCm Systems Profiler provides several additional features and utilities for causal profiling:
+
+.. csv-table::
+   :header: "Feature", "Coz", "ROCm Systems Profiler", "Notes"
+   :widths: 20, 60, 60, 30
+
+   "Debug info", "requires debug info in DWARF v3 format (``-gdwarf-3``)", "optional, supports any DWARF format version", "See Note #1 below"
+   "Experiment selection", "``<file>:<line>``", "``<function>`` or ``<file>:<line>``", "See Note #2 below"
+   "Experiment speed-ups", "Randomly samples b/t 0..100 in increments of 5 or one fixed speed-up", "Supports specifying smaller subset", "See Note #3 below"
+   "Scope options", "Supports binary and source scopes", "Supports binary, source, and function scopes", "See Note #4, #5, and #6 below"
+   "Scope inclusion", "Uses ``%`` as a wildcard for binary and source scopes", "Full regex support for binary, source, and function scopes", ""
+   "Scope exclusion", "Not supported", "Supports regexes for excluding binary/source/function", "See Note #7 below"
+   "Call-stack sampling", "Linux Perf", "Linux Perf, libunwind", "See Note #8 below"
+
+.. note::
+
+  #. ROCm Systems Profiler supports a "function" mode which does not require debug info.
+  #. ROCm Systems Profiler supports selecting an entire range of instruction pointers for a function instead
+     of an instruction pointer for one line. In large code bases, "function" mode
+     can resolve in fewer iterations. After a target function is identified, you can
+     switch to line mode and limit the function scope to the target function.
+  #. ROCm Systems Profiler supports randomly sampling from subsets, e.g. { 0, 0, 5, 10 }
+     where 0% is randomly selected 50% of time and 5% and 10% are randomly selected 25% of the time.
+  #. ROCm Systems Profiler and COZ have the same definition for binary scope, which is the binaries
+     loaded at runtime (the executable and linked libraries).
+  #. ROCm Systems Profiler "source scope" supports both ``<file>`` and ``<file>:<line>`` formats
+     in contrast to the COZ "source scope" which requires ``<file>:<line>`` format.
+  #. ROCm Systems Profiler supports a "function" scope which narrows the function and lines
+     which are eligible for causal experiments to those within the matching functions.
+  #. ROCm Systems Profiler supports a second filter on scopes for removing binary/source/function
+     caught by an inclusive match. For example ``BINARY_SCOPE=.*`` and ``BINARY_EXCLUDE=libmpi.*``
+     initially includes all binaries but exclude regex removes MPI libraries.
+  #. In ROCm Systems Profiler, the Linux Perf backend is preferred over use libunwind. However,
+     Linux Perf usage can be restricted for security reasons.
+     ROCm Systems Profiler falls back to using a second POSIX timer and libunwind if
+     Linux Perf is not available.

projects/rocprofiler-systems/docs/how-to/profiling-python-scripts.rst

@@ -0,0 +1,340 @@
+.. meta::
+   :description: ROCm Systems Profiler Python profiling documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, Python, profiling Python, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Profiling Python scripts
+****************************************************
+
+`ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ supports profiling Python code at the
+source level and the script level.
+Python support is enabled via the ``ROCPROFSYS_USE_PYTHON`` and the
+``ROCPROFSYS_PYTHON_VERSIONS="<MAJOR>.<MINOR>`` CMake options.
+Alternatively, to build multiple Python versions, use
+``ROCPROFSYS_PYTHON_VERSIONS="<MAJOR>.<MINOR>;[<MAJOR>.<MINOR>]"``,
+and ``ROCPROFSYS_PYTHON_ROOT_DIRS="/path/to/version;[/path/to/version]"`` instead of ``ROCPROFSYS_PYTHON_VERSION``.
+When building multiple Python versions, the length of the ``ROCPROFSYS_PYTHON_VERSIONS``
+and ``ROCPROFSYS_PYTHON_ROOT_DIRS`` lists must
+be the same size.
+
+.. note::
+
+   When using ROCm Systems Profiler with Python programs, the Python interpreter major and minor version (e.g. 3.7)
+   must match the interpreter major and minor version
+   used when compiling the Python bindings. When building ROCm Systems Profiler,
+   the shared object file ``libpyrocprofsys.<IMPL>-<VERSION>-<ARCH>-<OS>-<ABI>.so`` is generated
+   where ``IMPL`` is the Python implementation, ``VERSION`` is the major and minor
+   version, ``ARCH`` is the architecture,
+   ``OS`` is the operating system, and ``ABI`` is the application binary interface,
+   for example, ``libpyrocprofsys.cpython-38-x86_64-linux-gnu.so``.
+
+.. note::
+
+   ROCm Systems Profiler has limited support for Artificial Intelligence (AI) and Machine Learning (ML) workloads.
+   Data from child threads is not captured. For other profiling options,
+   see `rocprofV3 <https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/how-to/using-rocprofv3.html#using-rocprofv3>`_.
+
+Getting started
+========================================
+
+The ROCm Systems Profiler Python package is installed in ``lib/pythonX.Y/site-packages/rocprofsys``.
+To ensure the Python interpreter can find the ROCm Systems Profiler package,
+add this path to the ``PYTHONPATH`` environment variable, as in the following example:
+
+.. code-block:: shell
+
+   export PYTHONPATH=/opt/rocprofiler-systems/lib/python3.8/site-packages:${PYTHONPATH}
+
+Both the ``share/rocprofiler-systems/setup-env.sh`` script and the module file in
+``share/modulefiles/rocprofiler-systems`` automatically handle the prefixing of the ``PYTHONPATH``
+environment variable.
+
+Running ROCm Systems Profiler on a Python script
+================================================
+
+ROCm Systems Profiler provides an ``rocprof-sys-python`` helper bash script which
+ensures ``PYTHONPATH`` is properly set and the correct Python interpreter is used.
+This means the following commands are effectively equivalent:
+
+.. code-block:: shell
+
+   rocprof-sys-python --help
+
+and
+
+.. code-block:: shell
+
+   export PYTHONPATH=/opt/rocprofiler-systems/lib/python3.8/site-packages:${PYTHONPATH}
+   python3.8 -m rocprofsys --help
+
+.. note::
+
+   ``rocprof-sys-python`` and ``python -m rocprofsys`` use the same command-line syntax
+   as the other ``rocprof-sys`` executables (``rocprof-sys-python <ROCPROFSYS_ARGS> -- <SCRIPT> <SCRIPT_ARGS>``)
+   and has similar options.
+
+Command line options
+-----------------------------------
+
+Use ``rocprof-sys-python --help`` to view the available options:
+
+.. code-block:: shell
+
+   usage: rocprof-sys [-h] [-v VERBOSITY] [-b] [-c FILE] [-s FILE] [-F [BOOL]] [--label [{args,file,line} [{args,file,line} ...]]] [-I FUNC [FUNC ...]] [-E FUNC [FUNC ...]] [-R FUNC [FUNC ...]] [-MI FILE [FILE ...]] [-ME FILE [FILE ...]] [-MR FILE [FILE ...]] [--trace-c [BOOL]]
+
+   optional arguments:
+   -h, --help            show this help message and exit
+   -v VERBOSITY, --verbosity VERBOSITY
+                           Logging verbosity
+   -b, --builtin         Put 'profile' in the builtins. Use '@profile' to decorate a single function, or 'with profile:' to profile a single section of code.
+   -c FILE, --config FILE
+                           ROCm Systems Profiler configuration file
+   -s FILE, --setup FILE
+                           Code to execute before the code to profile
+   -F [BOOL], --full-filepath [BOOL]
+                           Encode the full function filename (instead of basename)
+   --label [{args,file,line} [{args,file,line} ...]]
+                           Encode the function arguments, filename, and/or line number into the profiling function label
+   -I FUNC [FUNC ...], --function-include FUNC [FUNC ...]
+                           Include any entries with these function names
+   -E FUNC [FUNC ...], --function-exclude FUNC [FUNC ...]
+                           Filter out any entries with these function names
+   -R FUNC [FUNC ...], --function-restrict FUNC [FUNC ...]
+                           Select only entries with these function names
+   -MI FILE [FILE ...], --module-include FILE [FILE ...]
+                           Include any entries from these files
+   -ME FILE [FILE ...], --module-exclude FILE [FILE ...]
+                           Filter out any entries from these files
+   -MR FILE [FILE ...], --module-restrict FILE [FILE ...]
+                           Select only entries from these files
+   --trace-c [BOOL]      Enable profiling C functions
+
+   usage: python3 -m rocprofsys <ROCPROFSYS_ARGS> -- <SCRIPT> <SCRIPT_ARGS>
+
+.. note::
+
+   The ``--trace-c`` option does not incorporate ROCm Systems Profiler's dynamic instrumentation support.
+   It only enables profiling the underlying C function call within the Python interpreter.
+
+Selective instrumentation
+-----------------------------------
+
+Similar to the ``rocprof-sys-instrument`` executable, command-line options exist for restricting,
+including, and excluding certain functions and modules, for example, ``--function-exclude "^__init__$"``.
+Alternatively, add the ``@profile`` decorator to the primary function of interest
+in your program and use the ``-b`` / ``--builtin`` command-line option to narrow the scope of the
+instrumentation to this function and its children.
+
+Consider the following Python code (``example.py``):
+
+.. code-block:: python
+
+   import sys
+
+   def fib(n):
+      return n if n < 2 else (fib(n - 1) + fib(n - 2))
+
+
+   def inefficient(n):
+      a = 0
+      for i in range(n):
+         a += i
+         for j in range(n):
+               a += j
+      return a
+
+
+   def run(n):
+      return fib(n) + inefficient(n)
+
+
+   if __name__ == "__main__":
+      run(20)
+
+Running ``rocprof-sys-python -- ./example.py`` with ``ROCPROFSYS_PROFILE=ON`` and
+``ROCPROFSYS_TIMEMORY_COMPONENTS=trip_count`` produces the following:
+
+.. code-block:: shell
+
+   |-------------------------------------------------------------------------------------------|
+   |                                COUNTS NUMBER OF INVOCATIONS                               |
+   |-------------------------------------------------------------------------------------------|
+   |                      LABEL                        | COUNT  | DEPTH  |   METRIC   |  SUM   |
+   |---------------------------------------------------|--------|--------|------------|--------|
+   | |0>>> run                                         |      1 |      0 | trip_count |      1 |
+   | |0>>> |_fib                                       |      1 |      1 | trip_count |      1 |
+   | |0>>>   |_fib                                     |      2 |      2 | trip_count |      2 |
+   | |0>>>     |_fib                                   |      4 |      3 | trip_count |      4 |
+   | |0>>>       |_fib                                 |      8 |      4 | trip_count |      8 |
+   | |0>>>         |_fib                               |     16 |      5 | trip_count |     16 |
+   | |0>>>           |_fib                             |     32 |      6 | trip_count |     32 |
+   | |0>>>             |_fib                           |     64 |      7 | trip_count |     64 |
+   | |0>>>               |_fib                         |    128 |      8 | trip_count |    128 |
+   | |0>>>                 |_fib                       |    256 |      9 | trip_count |    256 |
+   | |0>>>                   |_fib                     |    512 |     10 | trip_count |    512 |
+   | |0>>>                     |_fib                   |   1024 |     11 | trip_count |   1024 |
+   | |0>>>                       |_fib                 |   2026 |     12 | trip_count |   2026 |
+   | |0>>>                         |_fib               |   3632 |     13 | trip_count |   3632 |
+   | |0>>>                           |_fib             |   5020 |     14 | trip_count |   5020 |
+   | |0>>>                             |_fib           |   4760 |     15 | trip_count |   4760 |
+   | |0>>>                               |_fib         |   2942 |     16 | trip_count |   2942 |
+   | |0>>>                                 |_fib       |   1152 |     17 | trip_count |   1152 |
+   | |0>>>                                   |_fib     |    274 |     18 | trip_count |    274 |
+   | |0>>>                                     |_fib   |     36 |     19 | trip_count |     36 |
+   | |0>>>                                       |_fib |      2 |     20 | trip_count |      2 |
+   | |0>>> |_inefficient                               |      1 |      1 | trip_count |      1 |
+   |-------------------------------------------------------------------------------------------|
+
+If the ``inefficient`` function is decorated with ``@profile`` as follows:
+
+.. code-block:: python
+
+   @profile
+   def inefficient(n):
+      # ...
+
+And then run using the command ``rocprof-sys-python -b -- ./example.py``, ROCm Systems Profiler produces this output:
+
+.. code-block:: shell
+
+   |-----------------------------------------------------------|
+   |                COUNTS NUMBER OF INVOCATIONS               |
+   |-----------------------------------------------------------|
+   |      LABEL        | COUNT  | DEPTH  |   METRIC   |  SUM   |
+   |-------------------|--------|--------|------------|--------|
+   | |0>>> inefficient |      1 |      0 | trip_count |      1 |
+   |-----------------------------------------------------------|
+
+ROCm Systems Profiler Python source instrumentation
+===================================================
+
+Starting with the unmodified ``example.py`` script above, import the ``rocprofsys`` module:
+
+.. code-block:: python
+
+   import sys
+   import rocprofsys  # import rocprofsys
+
+   def fib(n):
+      # ... etc. ...
+
+Next, add ``@rocprofsys.profile()`` to the ``run`` function:
+
+.. code-block:: python
+
+   @rocprofsys.profile()
+   def run(n):
+      # ...
+
+Alternatively, use ``rocprofsys.profile()`` as a context-manager around ``run(20)``:
+
+.. code-block:: python
+
+   if __name__ == "__main__":
+      with rocprofsys.profile():
+         run(20)
+
+The results for both of the source-level instrumentation modes are identical to the
+original ``rocprof-sys-python -- ./example.py`` results:
+
+.. code-block:: shell
+
+   |-------------------------------------------------------------------------------------------|
+   |                                COUNTS NUMBER OF INVOCATIONS                               |
+   |-------------------------------------------------------------------------------------------|
+   |                      LABEL                        | COUNT  | DEPTH  |   METRIC   |  SUM   |
+   |---------------------------------------------------|--------|--------|------------|--------|
+   | |0>>> run                                         |      1 |      0 | trip_count |      1 |
+   | |0>>> |_fib                                       |      1 |      1 | trip_count |      1 |
+   | |0>>>   |_fib                                     |      2 |      2 | trip_count |      2 |
+   | |0>>>     |_fib                                   |      4 |      3 | trip_count |      4 |
+   | |0>>>       |_fib                                 |      8 |      4 | trip_count |      8 |
+   | |0>>>         |_fib                               |     16 |      5 | trip_count |     16 |
+   | |0>>>           |_fib                             |     32 |      6 | trip_count |     32 |
+   | |0>>>             |_fib                           |     64 |      7 | trip_count |     64 |
+   | |0>>>               |_fib                         |    128 |      8 | trip_count |    128 |
+   | |0>>>                 |_fib                       |    256 |      9 | trip_count |    256 |
+   | |0>>>                   |_fib                     |    512 |     10 | trip_count |    512 |
+   | |0>>>                     |_fib                   |   1024 |     11 | trip_count |   1024 |
+   | |0>>>                       |_fib                 |   2026 |     12 | trip_count |   2026 |
+   | |0>>>                         |_fib               |   3632 |     13 | trip_count |   3632 |
+   | |0>>>                           |_fib             |   5020 |     14 | trip_count |   5020 |
+   | |0>>>                             |_fib           |   4760 |     15 | trip_count |   4760 |
+   | |0>>>                               |_fib         |   2942 |     16 | trip_count |   2942 |
+   | |0>>>                                 |_fib       |   1152 |     17 | trip_count |   1152 |
+   | |0>>>                                   |_fib     |    274 |     18 | trip_count |    274 |
+   | |0>>>                                     |_fib   |     36 |     19 | trip_count |     36 |
+   | |0>>>                                       |_fib |      2 |     20 | trip_count |      2 |
+   | |0>>> |_inefficient                               |      1 |      1 | trip_count |      1 |
+   |-------------------------------------------------------------------------------------------|
+
+.. note::
+
+   When ``rocprof-sys-python`` is used without built-ins, the profiling results can be cluttered by the
+   numerous functions called when more complex modules are imported, such as ``import numpy``.
+
+ROCm Systems Profiler Python source instrumentation configuration
+-----------------------------------------------------------------
+
+Within the Python source code, the profiler can be configured by directly
+modifying the ``rocprof-sys.profiler.config`` data fields.
+
+.. code-block:: python
+
+   import sys
+
+   def fib(n):
+      return n if n < 2 else (fib(n - 1) + fib(n - 2))
+
+
+   def inefficient(n):
+      a = 0
+      for i in range(n):
+         a += i
+         for j in range(n):
+               a += j
+      return a
+
+
+   def run(n):
+      return fib(n) + inefficient(n)
+
+
+   if __name__ == "__main__":
+      from rocprofsys.profiler import config
+      from rocprofsys import profile
+
+      config.include_args = True
+      config.include_filename = False
+      config.include_line = False
+      config.restrict_functions += ["fib", "run"]
+
+      with profile():
+         run(5)
+
+Executing this script produces the following:
+
+.. code-block:: shell
+
+   |------------------------------------------------------------------|
+   |                   COUNTS NUMBER OF INVOCATIONS                   |
+   |------------------------------------------------------------------|
+   |          LABEL           | COUNT  | DEPTH  |   METRIC   |  SUM   |
+   |--------------------------|--------|--------|------------|--------|
+   | |0>>> run(n=5)           |      1 |      0 | trip_count |      1 |
+   | |0>>> |_fib(n=5)         |      1 |      1 | trip_count |      1 |
+   | |0>>>   |_fib(n=4)       |      1 |      2 | trip_count |      1 |
+   | |0>>>     |_fib(n=3)     |      1 |      3 | trip_count |      1 |
+   | |0>>>       |_fib(n=2)   |      1 |      4 | trip_count |      1 |
+   | |0>>>         |_fib(n=1) |      1 |      5 | trip_count |      1 |
+   | |0>>>         |_fib(n=0) |      1 |      5 | trip_count |      1 |
+   | |0>>>       |_fib(n=1)   |      1 |      4 | trip_count |      1 |
+   | |0>>>     |_fib(n=2)     |      1 |      3 | trip_count |      1 |
+   | |0>>>       |_fib(n=1)   |      1 |      4 | trip_count |      1 |
+   | |0>>>       |_fib(n=0)   |      1 |      4 | trip_count |      1 |
+   | |0>>>   |_fib(n=3)       |      1 |      2 | trip_count |      1 |
+   | |0>>>     |_fib(n=2)     |      1 |      3 | trip_count |      1 |
+   | |0>>>       |_fib(n=1)   |      1 |      4 | trip_count |      1 |
+   | |0>>>       |_fib(n=0)   |      1 |      4 | trip_count |      1 |
+   | |0>>>     |_fib(n=1)     |      1 |      3 | trip_count |      1 |
+   |------------------------------------------------------------------|

projects/rocprofiler-systems/docs/how-to/sampling-call-stack.rst

@@ -0,0 +1,413 @@
+.. meta::
+   :description: ROCm Systems Profiler call stack sampling documentation and reference
+   :keywords: rocprofiler-systems,rocprofsys, ROCm, profiler, sampling, call stack, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Sampling the call stack
+****************************************************
+
+`ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ can use call-stack sampling
+on a binary instrumented with either the ``rocprof-sys`` executable
+or the ``rocprof-sys-sample`` executable.
+For example, all of the following commands are effectively equivalent:
+
+* Binary rewrite with only the instrumentation necessary to start and stop sampling
+
+  .. code-block:: shell
+
+     rocprof-sys-instrument -M sampling -o foo.inst -- foo
+     rocprof-sys-run -- ./foo.inst
+
+* Runtime instrumentation with only the instrumentation necessary to start and stop sampling
+
+  .. code-block:: shell
+
+     rocprof-sys-instrument -M sampling -- foo
+
+* No instrumentation required
+
+  .. code-block:: shell
+
+     rocprof-sys-sample -- foo
+
+.. note::
+
+   Set ``ROCPROFSYS_USE_SAMPLING=ON`` to activate call-stack sampling when executing an instrumented binary.
+
+All ``rocprof-sys-instrument -M sampling`` (subsequently referred to as "instrumented-sampling")
+does is wrap the ``main`` of the executable with initialization
+before ``main`` starts and finalization after ``main`` ends.
+This can be accomplished without instrumentation through a ``LD_PRELOAD``
+of a library containing a dynamic symbol wrapper around ``__libc_start_main``.
+
+The use of ``rocprof-sys-sample`` is **recommended** over
+``rocprof-sys-instrument -M sampling`` when binary instrumentation
+is not necessary. This is for a number of reasons:
+
+* ``rocprof-sys-sample`` provides command-line options for controlling the ROCm Systems Profiler feature set instead of
+  requiring configuration files or environment variables
+* Despite the fact that instrumented-sampling only requires inserting snippets
+  around one function (``main``), Dyninst
+  does not have a feature for specifying that parsing and processing all the
+  other symbols in the binary is unnecessary.
+  In the best-case scenario when the target binary is relatively small,
+  instrumented-sampling has a slightly slower launch time,
+  but in the worst case scenarios it requires a significant amount of time and memory to launch.
+* ``rocprof-sys-sample`` is fully compatible with MPI. For example,
+  the command ``mpirun -n 2 rocprof-sys-sample -- foo`` is valid,
+  whereas ``mpirun -n 2 rocprof-sys-instrument -M sampling -- foo``
+  is incompatible with some MPI distributions (particularly OpenMPI). This is because
+  MPI prohibits forking within an MPI rank.
+
+  * When MPI and binary instrumentation are both involved, two steps are required:
+    performing a binary rewrite of the executable and then using the instrumented executable
+    in lieu of the original executable. ``rocprof-sys-sample`` is therefore much easier to use with MPI.
+
+The rocprof-sys-sample executable
+========================================
+
+View the help menu of ``rocprof-sys-sample`` with the ``-h`` / ``--help`` option:
+
+.. code-block:: shell
+
+   $ rocprof-sys-sample --help
+   [rocprof-sys-sample] Usage: rocprof-sys-sample [ --help (count: 0, dtype: bool)
+                                                --version (count: 0, dtype: bool)
+                                                --monochrome (max: 1, dtype: bool)
+                                                --debug (max: 1, dtype: bool)
+                                                --verbose (count: 1)
+                                                --config (min: 0, dtype: filepath)
+                                                --output (min: 1)
+                                                --trace (max: 1, dtype: bool)
+                                                --profile (max: 1, dtype: bool)
+                                                --flat-profile (max: 1, dtype: bool)
+                                                --host (max: 1, dtype: bool)
+                                                --device (max: 1, dtype: bool)
+                                                --wait (count: 1)
+                                                --duration (count: 1)
+                                                --trace-file (count: 1, dtype: filepath)
+                                                --trace-buffer-size (count: 1, dtype: KB)
+                                                --trace-fill-policy (count: 1)
+                                                --trace-wait (count: 1)
+                                                --trace-duration (count: 1)
+                                                --trace-periods (min: 1)
+                                                --trace-clock-id (count: 1)
+                                                --profile-format (min: 1)
+                                                --profile-diff (min: 1)
+                                                --process-freq (count: 1)
+                                                --process-wait (count: 1)
+                                                --process-duration (count: 1)
+                                                --cpus (count: unlimited, dtype: int or range)
+                                                --gpus (count: unlimited, dtype: int or range)
+                                                --freq (count: 1)
+                                                --sampling-wait (count: 1)
+                                                --sampling-duration (count: 1)
+                                                --tids (min: 1)
+                                                --cputime (min: 0)
+                                                --realtime (min: 0)
+                                                --include (count: unlimited)
+                                                --exclude (count: unlimited)
+                                                --cpu-events (count: unlimited)
+                                                --gpu-events (count: unlimited)
+                                                --inlines (max: 1, dtype: bool)
+                                                --hsa-interrupt (count: 1, dtype: int)
+                                             ]
+   Options:
+      -h, -?, --help                 Shows this page (count: 0, dtype: bool)
+      --version                      Prints the version and exit (count: 0, dtype: bool)
+
+      [DEBUG OPTIONS]
+
+      --monochrome                   Disable colorized output (max: 1, dtype: bool)
+      --debug                        Debug output (max: 1, dtype: bool)
+      -v, --verbose                  Verbose output (count: 1)
+
+      [GENERAL OPTIONS]  These are options which are ubiquitously applied
+
+      -c, --config                   Configuration file (min: 0, dtype: filepath)
+      -o, --output                   Output path. Accepts 1-2 parameters corresponding to the output path and the output prefix (min: 1)
+      -T, --trace                    Generate a detailed trace (perfetto output) (max: 1, dtype: bool)
+      -P, --profile                  Generate a call-stack-based profile (conflicts with --flat-profile) (max: 1, dtype: bool)
+      -F, --flat-profile             Generate a flat profile (conflicts with --profile) (max: 1, dtype: bool)
+      -H, --host                     Enable sampling host-based metrics for the process. E.g. CPU frequency, memory usage, etc. (max: 1, dtype: bool)
+      -D, --device                   Enable sampling device-based metrics for the process. E.g. GPU temperature, memory usage, etc. (max: 1, dtype: bool)
+      -w, --wait                     This option is a combination of '--trace-wait' and '--sampling-wait'. See the descriptions for those two options.
+                                    (count: 1)
+      -d, --duration                 This option is a combination of '--trace-duration' and '--sampling-duration'. See the descriptions for those two
+                                    options. (count: 1)
+
+      [TRACING OPTIONS]  Specific options controlling tracing (i.e. deterministic measurements of every event)
+
+      --trace-file                   Specify the trace output filename. Relative filepath will be with respect to output path and output prefix. (count: 1,
+                                    dtype: filepath)
+      --trace-buffer-size            Size limit for the trace output (in KB) (count: 1, dtype: KB)
+      --trace-fill-policy [ discard | ring_buffer ]
+
+                                    Policy for new data when the buffer size limit is reached:
+                                          - discard     : new data is ignored
+                                          - ring_buffer : new data overwrites oldest data (count: 1)
+      --trace-wait                   Set the wait time (in seconds) before collecting trace and/or profiling data(in seconds). By default, the duration is
+                                    in seconds of realtime but that can changed via --trace-clock-id. (count: 1)
+      --trace-duration               Set the duration of the trace and/or profile data collection (in seconds). By default, the duration is in seconds of
+                                    realtime but that can changed via --trace-clock-id. (count: 1)
+      --trace-periods                More powerful version of specifying trace delay and/or duration. Format is one or more groups of: <DELAY>:<DURATION>,
+                                    <DELAY>:<DURATION>:<REPEAT>, and/or <DELAY>:<DURATION>:<REPEAT>:<CLOCK_ID>. (min: 1)
+      --trace-clock-id [ 0 (realtime|CLOCK_REALTIME)
+                        1 (monotonic|CLOCK_MONOTONIC)
+                        2 (cputime|CLOCK_PROCESS_CPUTIME_ID)
+                        4 (monotonic_raw|CLOCK_MONOTONIC_RAW)
+                        5 (realtime_coarse|CLOCK_REALTIME_COARSE)
+                        6 (monotonic_coarse|CLOCK_MONOTONIC_COARSE)
+                        7 (boottime|CLOCK_BOOTTIME) ]
+                                    Set the default clock ID for for trace delay/duration. Note: "cputime" is the *process* CPU time and might need to be
+                                    scaled based on the number of threads, i.e. 4 seconds of CPU-time for an application with 4 fully active threads would
+                                    equate to ~1 second of realtime. If this proves to be difficult to handle in practice, please file a feature request
+                                    for rocprof-sys to auto-scale based on the number of threads. (count: 1)
+
+      [PROFILE OPTIONS]  Specific options controlling profiling (i.e. deterministic measurements which are aggregated into a summary)
+
+      --profile-format [ console | json | text ]
+                                    Data formats for profiling results (min: 1)
+      --profile-diff                 Generate a diff output b/t the profile collected and an existing profile from another run Accepts 1-2 parameters
+                                    corresponding to the input path and the input prefix (min: 1)
+
+      [HOST/DEVICE (PROCESS SAMPLING) OPTIONS]
+                                    Process sampling is background measurements for resources available to the entire process. These samples are not tied
+                                    to specific lines/regions of code
+
+      --process-freq                 Set the default host/device sampling frequency (number of interrupts per second) (count: 1)
+      --process-wait                 Set the default wait time (i.e. delay) before taking first host/device sample (in seconds of realtime) (count: 1)
+      --process-duration             Set the duration of the host/device sampling (in seconds of realtime) (count: 1)
+      --cpus                         CPU IDs for frequency sampling. Supports integers and/or ranges (count: unlimited, dtype: int or range)
+      --gpus                         GPU IDs for SMI queries. Supports integers and/or ranges (count: unlimited, dtype: int or range)
+
+      [GENERAL SAMPLING OPTIONS] General options for timer-based sampling per-thread
+
+      -f, --freq                     Set the default sampling frequency (number of interrupts per second) (count: 1)
+      --sampling-wait                Set the default wait time (i.e. delay) before taking first sample (in seconds). This delay time is based on the clock
+                                    of the sampler, i.e., a delay of 1 second for CPU-clock sampler may not equal 1 second of realtime (count: 1)
+      --sampling-duration            Set the duration of the sampling (in seconds of realtime). I.e., it is possible (currently) to set a CPU-clock time
+                                    delay that exceeds the real-time duration... resulting in zero samples being taken (count: 1)
+      -t, --tids                     Specify the default thread IDs for sampling, where 0 (zero) is the main thread and each thread created by the target
+                                    application is assigned an atomically incrementing value. (min: 1)
+
+      [SAMPLING TIMER OPTIONS] These options determine the heuristic for deciding when to take a sample
+
+      --cputime                      Sample based on a CPU-clock timer (default). Accepts zero or more arguments:
+                                          0. Enables sampling based on CPU-clock timer.
+                                          1. Interrupts per second. E.g., 100 == sample every 10 milliseconds of CPU-time.
+                                          2. Delay (in seconds of CPU-clock time). I.e., how long each thread should wait before taking first sample.
+                                          3+ Thread IDs to target for sampling, starting at 0 (the main thread).
+                                             May be specified as index or range, e.g., '0 2-4' will be interpreted as:
+                                                sample the main thread (0), do not sample the first child thread but sample the 2nd, 3rd, and 4th child threads (min: 0)
+      --realtime                     Sample based on a real-clock timer. Accepts zero or more arguments:
+                                          0. Enables sampling based on real-clock timer.
+                                          1. Interrupts per second. E.g., 100 == sample every 10 milliseconds of realtime.
+                                          2. Delay (in seconds of real-clock time). I.e., how long each thread should wait before taking first sample.
+                                          3+ Thread IDs to target for sampling, starting at 0 (the main thread).
+                                             May be specified as index or range, e.g., '0 2-4' will be interpreted as:
+                                                sample the main thread (0), do not sample the first child thread but sample the 2nd, 3rd, and 4th child threads
+                                             When sampling with a real-clock timer, please note that enabling this will cause threads which are typically "idle"
+                                             to consume more resources since, while idle, the real-clock time increases (and therefore triggers taking samples)
+                                             whereas the CPU-clock time does not. (min: 0)
+
+      [BACKEND OPTIONS]  These options control region information captured w/o sampling or instrumentation
+
+      -I, --include [ all | kokkosp | mpip | mutex-locks | ompt | rcclp | amd-smi | rocprofiler-sdk | rw-locks | spin-locks ]
+                                    Include data from these backends (count: unlimited)
+      -E, --exclude [ all | kokkosp | mpip | mutex-locks | ompt | rcclp | amd-smi | rocprofiler-sdk | rw-locks | spin-locks ]
+                                    Exclude data from these backends (count: unlimited)
+
+      [HARDWARE COUNTER OPTIONS] See also: rocprof-sys-avail -H
+
+      -C, --cpu-events               Set the CPU hardware counter events to record (ref: `rocprof-sys-avail -H -c CPU`) (count: unlimited)
+      -G, --gpu-events               Set the GPU hardware counter events to record (ref: `rocprof-sys-avail -H -c GPU`) (count: unlimited)
+
+      [MISCELLANEOUS OPTIONS]
+
+      -i, --inlines                  Include inline info in output when available (max: 1, dtype: bool)
+      --hsa-interrupt [ 0 | 1 ]      Set the value of the HSA_ENABLE_INTERRUPT environment variable.
+                                       ROCm version 5.2 and older have a bug which will cause a deadlock if a sample is taken while waiting for the signal
+                                       that a kernel completed -- which happens when sampling with a real-clock timer. We require this option to be set to
+                                       when --realtime is specified to make users aware that, while this may fix the bug, it can have a negative impact on
+                                       performance.
+                                       Values:
+                                          0     avoid triggering the bug, potentially at the cost of reduced performance
+                                          1     do not modify how ROCm is notified about kernel completion (count: 1, dtype: int)
+
+The general syntax for separating ROCm Systems Profiler command-line arguments from the
+following application arguments
+is consistent with the LLVM style of using a stand-alone double hyphen (``--``).
+All arguments preceding the double hyphen
+are interpreted as belonging to ROCm Systems Profiler and all arguments following it
+are interpreted as the
+application and its arguments. The double hyphen is only necessary when passing
+command-line arguments to a target
+which also uses hyphens. For example, you can run ``rocprof-sys-sample ls``, but
+to run ``ls -la``, use ``rocprof-sys-sample -- ls -la``.
+
+:doc:`Configuring the ROCm Systems Profiler runtime options <./configuring-runtime-options>`
+establishes the precedence of environment variable values over values specified
+in the configuration files. This enables
+you to configure the ROCm Systems Profiler runtime to your preferred default behavior
+in a file such as ``~/.rocprof-sys.cfg`` and then easily override
+those settings in the command line, for example, ``ROCPROFSYS_ENABLED=OFF rocprof-sys-sample -- foo``.
+Similarly, the command-line arguments passed to ``rocprof-sys-sample`` take precedence
+over environment variables.
+
+All of the command-line options above correlate to one or more configuration
+settings, for example, ``--cpu-events`` correlates to the ``ROCPROFSYS_PAPI_EVENTS`` configuration variable.
+``rocprof-sys-sample`` processes the arguments and outputs a summary of its configuration
+before running the target application.
+
+The following snippets show how ``rocprof-sys-sample`` runs with various environment updates.
+
+*  This snippet shows the environment updates when ``rocprof-sys-sample`` is invoked with no arguments:
+
+   .. code-block:: shell
+
+      $ rocprof-sys-sample -- ./parallel-overhead-locks 30 4 100
+
+      LD_PRELOAD=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCPROFSYS_USE_PROCESS_SAMPLING=false
+      ROCPROFSYS_USE_SAMPLING=true
+      OMP_TOOL_LIBRARIES=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCP_TOOL_LIB=/opt/rocprofiler-systems/lib/librocprof-sys.so.1.7.1
+
+*  The next snippet shows the environment updates when ``rocprof-sys-sample`` enables
+   profiling, tracing, device process-sampling, and does not enable host process-sampling:
+
+   .. code-block:: shell
+
+      $ rocprof-sys-sample -PTD -- ./parallel-overhead-locks 30 4 100
+
+      LD_PRELOAD=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCPROFSYS_CPU_FREQ_ENABLED=false
+      ROCPROFSYS_PROFILE=true
+      ROCPROFSYS_TRACE=true
+      ROCPROFSYS_USE_AMD_SMI=true
+      ROCPROFSYS_USE_PROCESS_SAMPLING=true
+      ROCPROFSYS_USE_SAMPLING=true
+      OMP_TOOL_LIBRARIES=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCP_TOOL_LIB=/opt/rocprofiler-systems/lib/librocprof-sys.so.1.7.1
+
+*  The next snippet shows the environment updates when ``rocprof-sys-sample`` enables
+   profiling, tracing, device process-sampling, host process-sampling, and all the available backends:
+
+   .. code-block:: shell
+
+      $ rocprof-sys-sample -PTDH -I all -- ./parallel-overhead-locks 30 4 100
+
+      KOKKOS_TOOLS_LIBS=/opt/rocprofiler-systems/lib/librocprof-sys.so.1.7.1
+      LD_PRELOAD=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCPROFSYS_CPU_FREQ_ENABLED=true
+      ROCPROFSYS_TRACE_THREAD_LOCKS=true
+      ROCPROFSYS_TRACE_THREAD_RW_LOCKS=true
+      ROCPROFSYS_TRACE_THREAD_SPIN_LOCKS=true
+      ROCPROFSYS_USE_KOKKOSP=true
+      ROCPROFSYS_USE_MPIP=true
+      ROCPROFSYS_USE_OMPT=true
+      ROCPROFSYS_TRACE=true
+      ROCPROFSYS_USE_PROCESS_SAMPLING=true
+      ROCPROFSYS_USE_RCCLP=true
+      ROCPROFSYS_USE_AMD_SMI=true
+      ROCPROFSYS_USE_ROCM=true
+      ROCPROFSYS_USE_SAMPLING=true
+      ROCPROFSYS_PROFILE=true
+      OMP_TOOL_LIBRARIES=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCP_TOOL_LIB=/opt/rocprofiler-systems/lib/librocprof-sys.so.1.7.1
+      ...
+
+*  The final snippet shows the environment updates when ``rocprof-sys-sample`` enables
+   profiling, tracing, device process-sampling, and host process-sampling,
+   sets the output path to ``rocprof-sys-output`` and the output prefix to ``%tag%``, and disables
+   all the available backends:
+
+   .. code-block:: shell
+
+      $ rocprof-sys-sample -PTDH -E all -o rocprof-sys-output %tag% -- ./parallel-overhead-locks 30 4 100
+
+      LD_PRELOAD=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.7.1
+      ROCPROFSYS_CPU_FREQ_ENABLED=true
+      ROCPROFSYS_OUTPUT_PATH=rocprof-sys-output
+      ROCPROFSYS_OUTPUT_PREFIX=%tag%
+      ROCPROFSYS_TRACE_THREAD_LOCKS=false
+      ROCPROFSYS_TRACE_THREAD_RW_LOCKS=false
+      ROCPROFSYS_TRACE_THREAD_SPIN_LOCKS=false
+      ROCPROFSYS_USE_KOKKOSP=false
+      ROCPROFSYS_USE_MPIP=false
+      ROCPROFSYS_USE_OMPT=false
+      ROCPROFSYS_TRACE=true
+      ROCPROFSYS_USE_PROCESS_SAMPLING=true
+      ROCPROFSYS_USE_RCCLP=false
+      ROCPROFSYS_USE_AMD_SMI=false
+      ROCPROFSYS_USE_ROCM=false
+      ROCPROFSYS_USE_SAMPLING=true
+      ROCPROFSYS_PROFILE=true
+      ...
+
+A rocprof-sys-sample example
+========================================
+
+Here is the full output from the previous
+``rocprof-sys-sample -PTDH -E all -o rocprof-sys-output %tag% -- ./parallel-overhead-locks 30 4 100`` command:
+
+.. code-block:: shell-session
+
+   $ rocprof-sys-sample -PTDH -E all -o rocprof-sys-output %tag% -c -- ./parallel-overhead-locks 30 4 100
+
+   LD_PRELOAD=/opt/rocprofiler-systems/lib/librocprof-sys-dl.so.1.11.3
+   ROCPROFSYS_CONFIG_FILE=
+   ROCPROFSYS_CPU_FREQ_ENABLED=true
+   ROCPROFSYS_OUTPUT_PATH=rocprof-sys-output
+   ROCPROFSYS_OUTPUT_PREFIX=%tag%
+   ROCPROFSYS_PROFILE=true
+   ROCPROFSYS_TRACE=true
+   ROCPROFSYS_TRACE_THREAD_LOCKS=false
+   ROCPROFSYS_TRACE_THREAD_RW_LOCKS=false
+   ROCPROFSYS_TRACE_THREAD_SPIN_LOCKS=false
+   ROCPROFSYS_USE_KOKKOSP=false
+   ROCPROFSYS_USE_MPIP=false
+   ROCPROFSYS_USE_OMPT=false
+   ROCPROFSYS_USE_PROCESS_SAMPLING=true
+   ROCPROFSYS_USE_RCCLP=false
+   ROCPROFSYS_USE_AMD_SMI=false
+   ROCPROFSYS_USE_ROCM=false
+   ROCPROFSYS_USE_SAMPLING=true
+   [rocprof-sys][dl][1785877] rocprofsys_main
+   [rocprof-sys][1785877][rocprofsys_init_tooling] Instrumentation mode: Sampling
+                                                     __
+       _ __    ___     ___   _ __    _ __    ___    / _|          ___   _   _   ___
+      | '__|  / _ \   / __| | '_ \  | '__|  / _ \  | |_   _____  / __| | | | | / __|
+      | |    | (_) | | (__  | |_) | | |    | (_) | |  _| |_____| \__ \ | |_| | \__ \
+      |_|     \___/   \___| | .__/  |_|     \___/  |_|           |___/  \__, | |___/
+                            |_|                                         |___/
+
+      rocprof-sys v1.11.2 (rev: 2586b74db8bf335742600010b8d9f1ce8da9cf89, compiler: GNU v11.4.1, rocm: v6.1.x)
+   [988.958]       perfetto.cc:58649 Configured tracing session 1, #sources:1, duration:0 ms, #buffers:1, total buffer size:1024000 KB, total sessions:1, uid:0 session name: ""
+   [parallel-overhead-locks] Threads: 4
+   [parallel-overhead-locks] Iterations: 100
+   [parallel-overhead-locks] fibonacci(30)...
+   [1] number of iterations: 100
+   [2] number of iterations: 100
+   [3] number of iterations: 100
+   [4] number of iterations: 100
+   [parallel-overhead-locks] fibonacci(30) x 4 = 409221992
+   [parallel-overhead-locks] number of mutex locks = 400
+   [rocprof-sys][1785877][0][rocprofsys_finalize] finalizing...
+   [rocprof-sys][1785877][0][rocprofsys_finalize]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] rocprof-sys/process/1785877 : 0.294342 sec wall_clock,    4.776 MB peak_rss,    3.170 MB page_rss, 0.990000 sec cpu_clock,  336.3 % cpu_util [laps: 1]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] rocprof-sys/process/1785877/thread/0 : 0.291535 sec wall_clock, 0.002619 sec thread_cpu_clock,    0.9 % thread_cpu_util,    4.776 MB peak_rss [laps: 1]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] rocprof-sys/process/1785877/thread/1 : 0.271353 sec wall_clock, 0.222572 sec thread_cpu_clock,   82.0 % thread_cpu_util,    4.200 MB peak_rss [laps: 1]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] rocprof-sys/process/1785877/thread/2 : 0.238218 sec wall_clock, 0.206405 sec thread_cpu_clock,   86.6 % thread_cpu_util,    3.432 MB peak_rss [laps: 1]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] rocprof-sys/process/1785877/thread/3 : 0.209459 sec wall_clock, 0.193415 sec thread_cpu_clock,   92.3 % thread_cpu_util,    2.472 MB peak_rss [laps: 1]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] rocprof-sys/process/1785877/thread/4 : 0.212029 sec wall_clock, 0.211694 sec thread_cpu_clock,   99.8 % thread_cpu_util,    1.152 MB peak_rss [laps: 1]
+   [rocprof-sys][1785877][0][rocprofsys_finalize]
+   [rocprof-sys][1785877][0][rocprofsys_finalize] Finalizing perfetto...
+   [rocprof-sys][1785877][perfetto]> Outputting '/home/user/code/rocprofiler-systems/build-release/rocprofiler-systems-output/2024-07-15_16.21/parallel-overhead-locksperfetto-trace-1785877.proto' (39.12 KB / 0.04 MB / 0.00 GB)... Done
+   [rocprof-sys][1785877][wall_clock]> Outputting 'rocprof-sys-output/2024-07-15_16.21/parallel-overhead-lockswall_clock-1785877.json'
+   [rocprof-sys][1785877][wall_clock]> Outputting 'rocprof-sys-output/2024-07-15_16.21/parallel-overhead-lockswall_clock-1785877.txt'
+   [rocprof-sys][1785877][metadata]> Outputting 'rocprof-sys-output/2024-07-15_16.21/parallel-overhead-locksmetadata-1785877.json' and 'rocprof-sys-output/2024-07-15_16.21/parallel-overhead-locksfunctions-1785877.json'
+   [rocprof-sys][1785877][0][rocprofsys_finalize] Finalized: 0.054582 sec wall_clock,    0.000 MB peak_rss,   -1.798 MB page_rss, 0.040000 sec cpu_clock,   73.3 % cpu_util
+   [989.312]       perfetto.cc:60128 Tracing session 1 ended, total sessions:0
+

projects/rocprofiler-systems/docs/how-to/using-rocprof-sys-api.rst

@@ -0,0 +1,313 @@
+.. meta::
+   :description: ROCm Systems Profiler documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Using the ROCm Systems Profiler API
+****************************************************
+
+The following example shows how a program can use the ROCm Systems Profiler API
+for run-time analysis.
+
+ROCm Systems Profiler user API example program
+==============================================
+
+You can use the ROCm Systems Profiler API to define custom regions to profile and trace.
+The following C++ program demonstrates this technique by calling several functions from the
+ROCm Systems Profiler API, such as ``rocprofsys_user_push_region`` and
+``rocprofsys_user_stop_thread_trace``.
+
+.. note::
+
+   By default, when ROCm Systems Profiler detects any ``rocprofsys_user_start_*`` or
+   ``rocprofsys_user_stop_*`` function, instrumentation
+   is disabled at start up, which means ``rocprofsys_user_stop_trace()`` is not
+   required at the beginning of ``main``. This behavior
+   can be manually controlled by using the ``ROCPROFSYS_INIT_ENABLED`` environment variable.
+   User-defined regions are always
+   recorded, regardless of whether ``rocprofsys_user_start_*`` or
+   ``rocprofsys_user_stop_*`` has been called.
+
+.. code-block:: shell
+
+   #include <rocprofiler-systems/categories.h>
+   #include <rocprofiler-systems/types.h>
+   #include <rocprofiler-systems/user.h>
+
+   #include <atomic>
+   #include <cassert>
+   #include <cerrno>
+   #include <cstdio>
+   #include <cstdlib>
+   #include <cstring>
+   #include <sstream>
+   #include <thread>
+   #include <vector>
+
+   std::atomic<long> total{ 0 };
+
+   long
+   fib(long n) __attribute__((noinline));
+
+   void
+   run(size_t nitr, long) __attribute__((noinline));
+
+   int
+   custom_push_region(const char* name);
+
+   namespace
+   {
+   rocprofsys_user_callbacks_t custom_callbacks   = ROCPROFSYS_USER_CALLBACKS_INIT;
+   rocprofsys_user_callbacks_t original_callbacks = ROCPROFSYS_USER_CALLBACKS_INIT;
+   }  // namespace
+
+   int
+   main(int argc, char** argv)
+   {
+      custom_callbacks.push_region = &custom_push_region;
+      rocprofsys_user_configure(ROCPROFSYS_USER_UNION_CONFIG, custom_callbacks,
+                              &original_callbacks);
+
+      rocprofsys_user_push_region(argv[0]);
+      rocprofsys_user_push_region("initialization");
+      size_t nthread = std::min<size_t>(16, std::thread::hardware_concurrency());
+      size_t nitr    = 50000;
+      long   nfib    = 10;
+      if(argc > 1) nfib = atol(argv[1]);
+      if(argc > 2) nthread = atol(argv[2]);
+      if(argc > 3) nitr = atol(argv[3]);
+      rocprofsys_user_pop_region("initialization");
+
+      printf("[%s] Threads: %zu\n[%s] Iterations: %zu\n[%s] fibonacci(%li)...\n", argv[0],
+            nthread, argv[0], nitr, argv[0], nfib);
+
+      rocprofsys_user_push_region("thread_creation");
+      std::vector<std::thread> threads{};
+      threads.reserve(nthread);
+      // disable instrumentation for child threads
+      rocprofsys_user_stop_thread_trace();
+      for(size_t i = 0; i < nthread; ++i)
+      {
+         threads.emplace_back(&run, nitr, nfib);
+      }
+      // re-enable instrumentation
+      rocprofsys_user_start_thread_trace();
+      rocprofsys_user_pop_region("thread_creation");
+
+      rocprofsys_user_push_region("thread_wait");
+      for(auto& itr : threads)
+         itr.join();
+      rocprofsys_user_pop_region("thread_wait");
+
+      run(nitr, nfib);
+
+      printf("[%s] fibonacci(%li) x %lu = %li\n", argv[0], nfib, nthread, total.load());
+      rocprofsys_user_pop_region(argv[0]);
+
+      return 0;
+   }
+
+   long
+   fib(long n)
+   {
+      return (n < 2) ? n : fib(n - 1) + fib(n - 2);
+   }
+
+   #define RUN_LABEL                                                                        \
+      std::string{ std::string{ __FUNCTION__ } + "(" + std::to_string(n) + ") x " +        \
+                  std::to_string(nitr) }                                                  \
+         .c_str()
+
+   void
+   run(size_t nitr, long n)
+   {
+      rocprofsys_user_push_region(RUN_LABEL);
+      long local = 0;
+      for(size_t i = 0; i < nitr; ++i)
+         local += fib(n);
+      total += local;
+      rocprofsys_user_pop_region(RUN_LABEL);
+   }
+
+   int
+   custom_push_region(const char* name)
+   {
+      if(!original_callbacks.push_region || !original_callbacks.push_annotated_region)
+         return ROCPROFSYS_USER_ERROR_NO_BINDING;
+
+      printf("Pushing custom region :: %s\n", name);
+
+      if(original_callbacks.push_annotated_region)
+      {
+         int32_t _err = errno;
+         char*   _msg = nullptr;
+         char    _buff[1024];
+         if(_err != 0) _msg = strerror_r(_err, _buff, sizeof(_buff));
+
+         rocprofsys_annotation_t _annotations[] = {
+               { "errno", ROCPROFSYS_INT32, &_err }, { "strerror", ROCPROFSYS_STRING, _msg }
+         };
+
+         errno = 0;  // reset errno
+         return (*original_callbacks.push_annotated_region)(
+               name, _annotations, sizeof(_annotations) / sizeof(rocprofsys_annotation_t));
+      }
+
+      return (*original_callbacks.push_region)(name);
+   }
+
+Linking the ROCm Systems Profiler libraries to another program
+==============================================================
+
+To link the ``rocprofiler-systems-user-library`` to another program,
+use the following CMake and ``g++`` directives.
+
+CMake
+-------------------------------------------------------
+
+.. code-block:: cmake
+
+   find_package(rocprofiler-systems REQUIRED COMPONENTS user)
+   add_executable(foo foo.cpp)
+   target_link_libraries(foo PRIVATE rocprofiler-systems::rocprofiler-systems-user-library)
+
+g++ compilation
+-------------------------------------------------------
+
+Assuming ROCm Systems Profiler is installed in ``/opt/rocprofsys``, use the ``g++`` compiler
+to build the application.
+
+.. code-block:: shell
+
+   g++ -g -I/opt/rocprofsys/include -L/opt/rocprofsys/lib foo.cpp -o foo -lrocprof-sys-user
+
+Output from the API example program
+========================================
+
+First, instrument and run the program.
+
+.. code-block:: shell-session
+
+   $ rocprof-sys-instrument -l --min-instructions=8 -E custom_push_region -o user-api.inst -- ./user-api
+   ...
+   $ rocprof-sys-run --profile --trace -- ./user-api.inst 10 12 1000
+
+   ROCPROFSYS: LD_PRELOAD=/home/rocm-dev/code/rocprofiler-systems/build/ubuntu/22.04/lib/librocprof-sys-dl.so.1.0.0
+   ROCPROFSYS: OMP_TOOL_LIBRARIES=/home/rocm-dev/code/rocprofiler-systems/build/ubuntu/22.04/lib/librocprof-sys-dl.so.1.0.0
+   ROCPROFSYS: ROCPROFSYS_PROFILE=true
+   ROCPROFSYS: ROCPROFSYS_TRACE=true
+   ROCPROFSYS: ROCPROFSYS_VERBOSE=0
+   [rocprof-sys][dl][1827155] rocprofsys_main
+   [rocprof-sys][1827155][rocprofsys_init_tooling] Instrumentation mode: Trace
+
+
+       ____   ___   ____ __  __   ______   ______ _____ _____ __  __ ____    ____  ____   ___  _____ ___ _     _____ ____
+      |  _ \ / _ \ / ___|  \/  | / ___\ \ / / ___|_   _| ____|  \/  / ___|  |  _ \|  _ \ / _ \|  ___|_ _| |   | ____|  _ \
+      | |_) | | | | |   | |\/| | \___ \\ V /\___ \ | | |  _| | |\/| \___ \  | |_) | |_) | | | | |_   | || |   |  _| | |_) |
+      |  _ <| |_| | |___| |  | |  ___) || |  ___) || | | |___| |  | |___) | |  __/|  _ <| |_| |  _|  | || |___| |___|  _ <
+      |_| \_\\___/ \____|_|  |_| |____/ |_| |____/ |_| |_____|_|  |_|____/  |_|   |_| \_\\___/|_|   |___|_____|_____|_| \_\
+
+      rocprof-sys v1.0.0 (rev: 3213dc652728f7ed01b62bf55f6af76c43bfcbdb, x86_64-linux-gnu, compiler: GNU v11.4.0, rocm: v6.3.x)
+   [790.763]       perfetto.cc:47606 Configured tracing session 1, #sources:1, duration:0 ms, #buffers:1, total buffer size:1024000 KB, total sessions:1, uid:0 session name: ""
+   [./user-api.inst] Threads: 12
+   [./user-api.inst] Iterations: 1000
+   [./user-api.inst] fibonacci(10)...
+   [./user-api.inst] fibonacci(10) x 12 = 715000
+
+   [rocprof-sys][1827155][0][rocprofsys_finalize] finalizing...
+   [rocprof-sys][1827155][0][rocprofsys_finalize]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155 : 0.137404 sec wall_clock,    6.528 MB peak_rss,    6.685 MB page_rss, 0.540000 sec cpu_clock,  393.0 % cpu_util [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/0 : 0.135815 sec wall_clock, 0.035171 sec thread_cpu_clock,   25.9 % thread_cpu_util,    6.016 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/1 : 0.028336 sec wall_clock, 0.028336 sec thread_cpu_clock,  100.0 % thread_cpu_util,    0.640 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/2 : 0.030380 sec wall_clock, 0.030380 sec thread_cpu_clock,  100.0 % thread_cpu_util,    3.840 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/3 : 0.035233 sec wall_clock, 0.035227 sec thread_cpu_clock,  100.0 % thread_cpu_util,    3.840 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/4 : 0.035275 sec wall_clock, 0.035267 sec thread_cpu_clock,  100.0 % thread_cpu_util,    3.840 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/5 : 0.035452 sec wall_clock, 0.035452 sec thread_cpu_clock,  100.0 % thread_cpu_util,    3.840 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/6 : 0.036198 sec wall_clock, 0.036190 sec thread_cpu_clock,  100.0 % thread_cpu_util,    3.840 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/7 : 0.034709 sec wall_clock, 0.034702 sec thread_cpu_clock,  100.0 % thread_cpu_util,    0.640 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/8 : 0.036590 sec wall_clock, 0.033590 sec thread_cpu_clock,   91.8 % thread_cpu_util,    0.512 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/9 : 0.033108 sec wall_clock, 0.033098 sec thread_cpu_clock,  100.0 % thread_cpu_util,    0.384 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/10 : 0.032993 sec wall_clock, 0.032994 sec thread_cpu_clock,  100.0 % thread_cpu_util,    0.256 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/11 : 0.035687 sec wall_clock, 0.035368 sec thread_cpu_clock,   99.1 % thread_cpu_util,    0.128 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] rocprofsys/process/1827155/thread/12 : 0.035359 sec wall_clock, 0.035354 sec thread_cpu_clock,  100.0 % thread_cpu_util,    0.128 MB peak_rss [laps: 1]
+   [rocprof-sys][1827155][0][rocprofsys_finalize]
+   [rocprof-sys][1827155][0][rocprofsys_finalize] Finalizing perfetto...
+   [rocprofiler-systems][1827155][perfetto]> Outputting '/home/rocm-dev/opt/user-api-test/rocprofsys-user-api.inst-output/2025-01-15_17.57/perfetto-trace-1827155.proto' (17.20 KB / 0.02 MB / 0.00 GB)... Done
+   [rocprofiler-systems][1827155][wall_clock]> Outputting 'rocprofsys-user-api.inst-output/2025-01-15_17.57/wall_clock-1827155.json'
+   [rocprofiler-systems][1827155][wall_clock]> Outputting 'rocprofsys-user-api.inst-output/2025-01-15_17.57/wall_clock-1827155.txt'
+   [rocprofiler-systems][1827155][metadata]> Outputting 'rocprofsys-user-api.inst-output/2025-01-15_17.57/metadata-1827155.json' and 'rocprofsys-user-api.inst-output/2025-01-15_17.57/functions-1827155.json'
+   [rocprof-sys][1827155][0][rocprofsys_finalize] Finalized: 0.048039 sec wall_clock,    0.640 MB peak_rss,    0.655 MB page_rss, 0.020000 sec cpu_clock,   41.6 % cpu_util
+   [790.953]       perfetto.cc:49204 Tracing session 1 ended, total sessions:0
+
+Then review the output.
+
+.. code-block:: shell
+
+   $ cat rocprof-sys-example-output/wall_clock.txt
+   |----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|
+   |                                                                                          REAL-CLOCK TIMER (I.E. WALL-CLOCK TIMER)                                                                                          |
+   |----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|
+   |                                                 LABEL                                                   | COUNT  | DEPTH  |   METRIC   | UNITS  |   SUM    |   MEAN   |   MIN    |   MAX    |   VAR    | STDDEV   | % SELF |
+   |---------------------------------------------------------------------------------------------------------|--------|--------|------------|--------|----------|----------|----------|----------|----------|----------|--------|
+   | |00>>> ./user-api.inst                                                                                  |      1 |      0 | wall_clock | sec    | 0.874293 | 0.874293 | 0.874293 | 0.874293 | 0.000000 | 0.000000 |    0.0 |
+   | |00>>> |_initialization                                                                                 |      1 |      1 | wall_clock | sec    | 0.000015 | 0.000015 | 0.000015 | 0.000015 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>> |_thread_creation                                                                                |      1 |      1 | wall_clock | sec    | 0.059934 | 0.059934 | 0.059934 | 0.059934 | 0.000000 | 0.000000 |    1.0 |
+   | |00>>>   |_pthread_create                                                                               |     12 |      2 | wall_clock | sec    | 0.059338 | 0.004945 | 0.004455 | 0.005743 | 0.000000 | 0.000479 |    0.0 |
+   | |01>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.027499 | 0.027499 | 0.027499 | 0.027499 | 0.000000 | 0.000000 |    0.1 |
+   | |01>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.027463 | 0.027463 | 0.027463 | 0.027463 | 0.000000 | 0.000000 |  100.0 |
+   | |02>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.027804 | 0.027804 | 0.027804 | 0.027804 | 0.000000 | 0.000000 |    0.2 |
+   | |02>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.027752 | 0.027752 | 0.027752 | 0.027752 | 0.000000 | 0.000000 |  100.0 |
+   | |03>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.027567 | 0.027567 | 0.027567 | 0.027567 | 0.000000 | 0.000000 |    0.1 |
+   | |03>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.027529 | 0.027529 | 0.027529 | 0.027529 | 0.000000 | 0.000000 |  100.0 |
+   | |05>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.027699 | 0.027699 | 0.027699 | 0.027699 | 0.000000 | 0.000000 |    0.2 |
+   | |05>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.027651 | 0.027651 | 0.027651 | 0.027651 | 0.000000 | 0.000000 |  100.0 |
+   | |04>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.033427 | 0.033427 | 0.033427 | 0.033427 | 0.000000 | 0.000000 |    0.2 |
+   | |04>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.033376 | 0.033376 | 0.033376 | 0.033376 | 0.000000 | 0.000000 |  100.0 |
+   | |06>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.032210 | 0.032210 | 0.032210 | 0.032210 | 0.000000 | 0.000000 |    0.1 |
+   | |06>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.032168 | 0.032168 | 0.032168 | 0.032168 | 0.000000 | 0.000000 |  100.0 |
+   | |07>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.030176 | 0.030176 | 0.030176 | 0.030176 | 0.000000 | 0.000000 |    0.2 |
+   | |07>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.030122 | 0.030122 | 0.030122 | 0.030122 | 0.000000 | 0.000000 |  100.0 |
+   | |08>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.027941 | 0.027941 | 0.027941 | 0.027941 | 0.000000 | 0.000000 |    0.1 |
+   | |08>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.027899 | 0.027899 | 0.027899 | 0.027899 | 0.000000 | 0.000000 |  100.0 |
+   | |09>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.034679 | 0.034679 | 0.034679 | 0.034679 | 0.000000 | 0.000000 |    0.1 |
+   | |09>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.034636 | 0.034636 | 0.034636 | 0.034636 | 0.000000 | 0.000000 |  100.0 |
+   | |11>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.028143 | 0.028143 | 0.028143 | 0.028143 | 0.000000 | 0.000000 |    0.1 |
+   | |11>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.028103 | 0.028103 | 0.028103 | 0.028103 | 0.000000 | 0.000000 |  100.0 |
+   | |10>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.033393 | 0.033393 | 0.033393 | 0.033393 | 0.000000 | 0.000000 |    0.1 |
+   | |10>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.033354 | 0.033354 | 0.033354 | 0.033354 | 0.000000 | 0.000000 |  100.0 |
+   | |12>>>     |_start_thread                                                                               |      1 |      3 | wall_clock | sec    | 0.027765 | 0.027765 | 0.027765 | 0.027765 | 0.000000 | 0.000000 |    0.2 |
+   | |12>>>       |_run(10) x 1000                                                                           |      1 |      4 | wall_clock | sec    | 0.027710 | 0.027710 | 0.027710 | 0.027710 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>> |_thread_wait                                                                                    |      1 |      1 | wall_clock | sec    | 0.027971 | 0.027971 | 0.027971 | 0.027971 | 0.000000 | 0.000000 |    1.3 |
+   | |00>>>   |_std::vector<std::thread, std::allocator<std::thread> >::begin                                |      1 |      2 | wall_clock | sec    | 0.000003 | 0.000003 | 0.000003 | 0.000003 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>   |_std::vector<std::thread, std::allocator<std::thread> >::end                                  |      1 |      2 | wall_clock | sec    | 0.000002 | 0.000002 | 0.000002 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>   |___gnu_cxx::operator!=<std::thread*, std::vector<std::thread, std::allocator<std::thread> > > |     13 |      2 | wall_clock | sec    | 0.000024 | 0.000002 | 0.000001 | 0.000003 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>   |_pthread_join                                                                                 |     12 |      2 | wall_clock | sec    | 0.027583 | 0.002299 | 0.000003 | 0.011250 | 0.000011 | 0.003289 |  100.0 |
+   | |00>>> |_run                                                                                            |      1 |      1 | wall_clock | sec    | 0.786236 | 0.786236 | 0.786236 | 0.786236 | 0.000000 | 0.000000 |    0.0 |
+   | |00>>>   |_std::char_traits<char>::length                                                               |      1 |      2 | wall_clock | sec    | 0.000002 | 0.000002 | 0.000002 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>   |_std::distance<char const*>                                                                   |      1 |      2 | wall_clock | sec    | 0.000002 | 0.000002 | 0.000002 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>   |_std::operator+<char, std::char_traits<char>, std::allocator<char> >                          |      4 |      2 | wall_clock | sec    | 0.000006 | 0.000002 | 0.000001 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>   |_run(10) x 1000                                                                               |      1 |      2 | wall_clock | sec    | 0.786184 | 0.786184 | 0.786184 | 0.786184 | 0.000000 | 0.000000 |    0.0 |
+   | |00>>>     |_run [{95,25}-{97,25}]                                                                      |      1 |      3 | wall_clock | sec    | 0.786141 | 0.786141 | 0.786141 | 0.786141 | 0.000000 | 0.000000 |    0.4 |
+   | |00>>>       |_fib                                                                                      |   1000 |      4 | wall_clock | sec    | 0.782692 | 0.000783 | 0.000757 | 0.001397 | 0.000000 | 0.000026 |    1.0 |
+   | |00>>>         |_fib                                                                                    |   2000 |      5 | wall_clock | sec    | 0.774875 | 0.000387 | 0.000282 | 0.000863 | 0.000000 | 0.000095 |    2.0 |
+   | |00>>>           |_fib                                                                                  |   4000 |      6 | wall_clock | sec    | 0.759351 | 0.000190 | 0.000101 | 0.000570 | 0.000000 | 0.000068 |    4.0 |
+   | |00>>>             |_fib                                                                                |   8000 |      7 | wall_clock | sec    | 0.728911 | 0.000091 | 0.000034 | 0.000350 | 0.000000 | 0.000042 |    8.5 |
+   | |00>>>               |_fib                                                                              |  16000 |      8 | wall_clock | sec    | 0.666793 | 0.000042 | 0.000009 | 0.000206 | 0.000000 | 0.000025 |   18.5 |
+   | |00>>>                 |_fib                                                                            |  32000 |      9 | wall_clock | sec    | 0.543524 | 0.000017 | 0.000001 | 0.000121 | 0.000000 | 0.000014 |   38.3 |
+   | |00>>>                   |_fib                                                                          |  52000 |     10 | wall_clock | sec    | 0.335118 | 0.000006 | 0.000001 | 0.000070 | 0.000000 | 0.000008 |   61.0 |
+   | |00>>>                     |_fib                                                                        |  44000 |     11 | wall_clock | sec    | 0.130629 | 0.000003 | 0.000001 | 0.000036 | 0.000000 | 0.000004 |   79.4 |
+   | |00>>>                       |_fib                                                                      |  16000 |     12 | wall_clock | sec    | 0.026893 | 0.000002 | 0.000001 | 0.000026 | 0.000000 | 0.000002 |   91.5 |
+   | |00>>>                         |_fib                                                                    |   2000 |     13 | wall_clock | sec    | 0.002279 | 0.000001 | 0.000001 | 0.000014 | 0.000000 | 0.000001 |  100.0 |
+   | |00>>>     |_std::char_traits<char>::length                                                             |      1 |      3 | wall_clock | sec    | 0.000001 | 0.000001 | 0.000001 | 0.000001 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>     |_std::distance<char const*>                                                                 |      1 |      3 | wall_clock | sec    | 0.000001 | 0.000001 | 0.000001 | 0.000001 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>>     |_std::operator+<char, std::char_traits<char>, std::allocator<char> >                        |      4 |      3 | wall_clock | sec    | 0.000005 | 0.000001 | 0.000001 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>> |_std::operator&                                                                                 |      1 |      1 | wall_clock | sec    | 0.000003 | 0.000003 | 0.000003 | 0.000003 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>> std::vector<std::thread, std::allocator<std::thread> >::~vector                                  |      1 |      0 | wall_clock | sec    | 0.000256 | 0.000256 | 0.000256 | 0.000256 | 0.000000 | 0.000000 |   20.9 |
+   | |00>>> |_std::thread::~thread                                                                           |     12 |      1 | wall_clock | sec    | 0.000193 | 0.000016 | 0.000014 | 0.000025 | 0.000000 | 0.000004 |   31.9 |
+   | |00>>>   |_std::thread::joinable                                                                        |     12 |      2 | wall_clock | sec    | 0.000131 | 0.000011 | 0.000010 | 0.000017 | 0.000000 | 0.000003 |   77.6 |
+   | |00>>>     |_std::thread::id::id                                                                        |     12 |      3 | wall_clock | sec    | 0.000016 | 0.000001 | 0.000001 | 0.000004 | 0.000000 | 0.000001 |  100.0 |
+   | |00>>>     |_std::operator==                                                                            |     12 |      3 | wall_clock | sec    | 0.000014 | 0.000001 | 0.000001 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>> |_std::allocator_traits<std::allocator<std::thread> >::deallocate                                |      1 |      1 | wall_clock | sec    | 0.000008 | 0.000008 | 0.000008 | 0.000008 | 0.000000 | 0.000000 |   70.8 |
+   | |00>>>   |___gnu_cxx::new_allocator<std::thread>::deallocate                                            |      1 |      2 | wall_clock | sec    | 0.000002 | 0.000002 | 0.000002 | 0.000002 | 0.000000 | 0.000000 |  100.0 |
+   | |00>>> |_std::allocator<std::thread>::~allocator                                                        |      1 |      1 | wall_clock | sec    | 0.000001 | 0.000001 | 0.000001 | 0.000001 | 0.000000 | 0.000000 |  100.0 |
+   |----------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------------|

projects/rocprofiler-systems/docs/how-to/vcn-jpeg-sampling.rst

@@ -0,0 +1,160 @@
+.. meta::
+   :description: ROCm Systems Profiler VCN and JPEG activity sampling and tracing
+   :keywords: rocprof-sys, rocprofiler-systems, ROCm, tips, how to, profiler, tracking, VCN, JPEG, rocDecode, rocjpeg, AMD
+
+********************************************
+VCN and JPEG activity sampling and tracing
+********************************************
+
+`ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ supports
+sampling of VCN and JPEG engines activities. It allows you to gather key performance metrics for
+VCN utilization and understand engine usage through visualization. This information can be used
+to optimize media and video workloads. Additionally, it supports tracing of `rocDecode
+<https://rocm.docs.amd.com/projects/rocDecode/en/latest/>`_ APIs, `rocJPEG
+<https://rocm.docs.amd.com/projects/rocJPEG/en/latest/>`_ APIs, and the
+Video Acceleration APIs (VA-APIs). Tracing these APIs provides insights into how different
+components of the video encoding and decoding workloads interact with the VCN engine.
+
+Sampling support
+=================
+
+Sampling of VCN and JPEG engine activity is supported by leveraging `AMD SMI <https://rocm.docs.amd.com/projects/amdsmi/en/latest/>`_ which provides the interface for GPU metric collection.
+
+1. Set the ``ROCPROFSYS_USE_AMD_SMI`` environment variable to enable GPU metric collection:
+
+.. code-block:: shell
+
+  export ROCPROFSYS_USE_AMD_SMI=true
+
+2. Update the ``ROCPROFSYS_AMD_SMI_METRICS`` variable to collect the VCN and JPEG activity metrics. The default value is:
+
+.. code-block:: shell
+
+  ROCPROFSYS_AMD_SMI_METRICS=busy,temp,power,mem_usage
+
+To include VCN and JPEG activity metrics, update it to:
+
+.. code-block:: shell
+
+  ROCPROFSYS_AMD_SMI_METRICS=busy,temp,power,mem_usage,vcn_activity,jpeg_activity
+
+Alternatively, you can use the following to collect all available GPU metrics:
+
+.. code-block:: shell
+
+  ROCPROFSYS_AMD_SMI_METRICS=all
+
+API tracing support
+=====================
+
+Tracing of rocDecode and rocJPEG APIs is supported by leveraging `ROCprofiler-SDK <https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/index.html>`_ 
+which provides runtime-independent APIs for tracing the runtime calls and asynchronous activities associated with decoder activities and workload in VCN and JPEG engines.
+
+To enable tracing for the rocDecode and rocJPEG APIs, update the ``ROCPROFSYS_ROCM_DOMAINS`` variable. The default value is:
+
+.. code-block:: shell
+
+  ROCPROFSYS_ROCM_DOMAINS=hip_runtime_api,marker_api,kernel_dispatch,memory_copy,scratch_memory,page_migration
+
+Add ``rocdecode_api`` and ``rocjpeg_api`` to include tracing for rocDecode and rocJPEG APIs:
+
+.. code-block:: shell
+
+  ROCPROFSYS_ROCM_DOMAINS=hip_runtime_api,marker_api,kernel_dispatch,memory_copy,scratch_memory,page_migration,rocdecode_api,rocjpeg_api
+
+.. note::
+
+   By default, enabling ``rocdecode_api`` or ``rocjpeg_api`` also enables VA-API tracing.
+
+To explore all supported tracing domains, use the command:
+
+.. code-block:: shell
+
+  rocprof-sys-avail -bd -r ROCM_DOMAINS
+
+For more details on the APIs, refer to `ROCprofiler-SDK Developer Docs <https://rocm.docs.amd.com/projects/rocprofiler-sdk/en/latest/_doxygen/rocprofiler-sdk/html/>`_.
+
+Using rocDecode and rocJPEG samples
+================================================
+
+For testing purposes, you can use the `rocDecode samples <https://github.com/ROCm/rocDecode?tab=readme-ov-file#using-sample-application>`_
+and `rocJPEG samples <https://github.com/ROCm/rocJPEG?tab=readme-ov-file#using-sample-application>`_.
+For generating sufficient load for VCN and JPEG engines, you can use the following samples:
+
+For video decoding:
+  - `Video decode batch <https://github.com/ROCm/rocDecode/tree/develop/samples/videoDecodeBatch>`_
+  - `Video decode performance <https://github.com/ROCm/rocDecode/tree/develop/samples/videoDecodePerf>`_
+
+For JPEG decoding:
+  - `JPEG decode batched <https://github.com/ROCm/rocJPEG/tree/develop/samples/jpegDecodeBatched>`_
+  - `JPEG decode perf <https://github.com/ROCm/rocJPEG/tree/develop/samples/jpegDecodePerf>`_
+
+After completing the build steps mentioned in the sample documentation, proceed with the following steps:
+
+1. Source the ROCm Systems Profiler Environment using:
+
+.. code-block:: shell
+
+   source /opt/rocprofiler-systems/share/rocprofiler-systems/setup-env.sh
+
+Alternatively, if you are using modules, use:
+
+.. code-block:: shell
+
+   module use /opt/rocprofiler-systems/share/modulefiles
+
+2. Generate and configure the profiler config file.
+
+.. code-block:: shell
+
+   rocprof-sys-avail -G $HOME/.rocprofsys.cfg -F txt
+   export ROCPROFSYS_CONFIG_FILE=$HOME/.rocprofsys.cfg
+
+Edit ``.rocprofsys.cfg`` with the following settings:
+
+.. code-block:: shell
+
+  ROCPROFSYS_USE_AMD_SMI     = true
+  ROCPROFSYS_AMD_SMI_METRICS = busy,temp,power,mem_usage,vcn_activity,jpeg_activity
+  ROCPROFSYS_ROCM_DOMAINS    = hip_runtime_api,marker_api,kernel_dispatch,memory_copy,scratch_memory,page_migration,rocdecode_api,rocjpeg_api
+
+3. Profile the rocDecode sample.
+
+.. code-block:: shell
+
+  rocprof-sys-sample -PTHD -- ./videodecodebatch -i /opt/rocm/share/rocdecode/video/
+
+.. note::
+
+   If the ``rocdecode-dev`` package is installed, then the sample videos will be located in ``/opt/rocm/share/rocdecode/video``, by default.
+
+At the end of the run, a similar message appears::
+
+  [rocprofiler-systems][964294][perfetto]> Outputting '/home/demo/rocprofsys-videodecodebatch-output/2025-04-25_15.52/perfetto-trace-964294.proto'
+  (2792.91 KB / 2.79 MB / 0.00 GB)... Done
+
+
+To view the generated ``.proto`` file in the browser, open the
+`Perfetto UI page <https://ui.perfetto.dev/>`_. Then, click on
+``Open trace file`` and select the ``.proto`` file. In the browser, a similar visualization is generated.
+
+.. image:: ../data/rocprof-sys-vcn-activity.png
+   :alt: Visualization of a performance graph in Perfetto with VCN Activity tracks
+
+.. image:: ../data/rocprof-sys-rocdecode.png
+   :alt: Visualization of a performance graph in Perfetto with rocdecode and VA-API traces
+
+4. To profile the rocJPEG sample, use:
+
+.. code-block:: shell
+
+  rocprof-sys-sample -v 2 -PTHD -- ./jpegdecodeperf -i /opt/rocm/share/rocjpeg/image/
+
+.. note::
+
+  If ``rocjpeg-dev`` package is installed, the sample images will be located in the
+  ``/opt/rocm/share/rocjpeg/image/`` directory.
+  Duplicate the images to generate enough workload to see activity in the trace
+
+.. image:: ../data/rocprof-sys-jpeg-activity.png
+   :alt: Visualization of a performance graph in Perfetto with JPEG Activity tracks

projects/rocprofiler-systems/docs/index.rst

@@ -0,0 +1,72 @@
+.. meta::
+   :description: ROCm Systems Profiler documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+***********************************
+ROCm Systems Profiler documentation
+***********************************
+
+ROCm Systems Profiler is designed for the high-level profiling and comprehensive tracing
+of applications running on the CPU or the CPU and GPU. It supports dynamic binary
+instrumentation, call-stack sampling, and various other features for determining
+which function and line number are currently executing. To learn more, see :doc:`what-is-rocprof-sys`
+
+ROCm Systems Profiler is open source and hosted at `<https://github.com/ROCm/rocprofiler-systems>`__.
+It is the successor to `<https://github.com/ROCm/omnitrace>`__.
+
+.. grid:: 2
+  :gutter: 3
+
+  .. grid-item-card:: Install
+
+    * :doc:`Quick start <./install/quick-start>`
+    * :doc:`ROCm Systems Profiler installation <./install/install>`
+
+Use the following topics to learn more about the advantages of ROCm Systems Profiler in application
+profiling, how it supports performance analysis, and how to leverage its capabilities in practice:
+
+.. grid:: 2
+  :gutter: 3
+
+  .. grid-item-card:: How to
+
+    * :doc:`Configuring the environment <./how-to/configuring-validating-environment>`
+
+      * :doc:`Configuring runtime options <./how-to/configuring-runtime-options>`
+
+    * :doc:`Profiling <./how-to/general-tips-using-rocprof-sys>`
+
+      * :doc:`Sampling the call stack <./how-to/sampling-call-stack>`
+      * :doc:`Instrumenting and rewriting a binary application <./how-to/instrumenting-rewriting-binary-application>`
+      * :doc:`Performing causal profiling <./how-to/performing-causal-profiling>`
+      * :doc:`Profiling Python scripts <./how-to/profiling-python-scripts>`
+      * :doc:`Network performance profiling <./how-to/nic-profiling>`
+      * :doc:`VCN and JPEG sampling and tracing <./how-to/vcn-jpeg-sampling>`
+
+    * :doc:`Understanding the output <./how-to/understanding-rocprof-sys-output>`
+    * :doc:`Using the ROCm Systems Profiler API <./how-to/using-rocprof-sys-api>`
+
+  .. grid-item-card:: Conceptual
+
+    * :doc:`Data collection modes <./conceptual/data-collection-modes>`
+    * :doc:`Features and use cases <./conceptual/rocprof-sys-feature-set>`
+
+  .. grid-item-card:: Reference
+
+    * :doc:`Development guide <./reference/development-guide>`
+    * :doc:`Glossary <./reference/rocprof-sys-glossary>`
+    * :doc:`API library <./doxygen/html/files>`
+    * :doc:`Class member functions <./doxygen/html/functions>`
+    * :doc:`Globals <./doxygen/html/globals>`
+    * :doc:`Classes, structures, and interfaces <./doxygen/html/annotated>`
+
+  .. grid-item-card:: Tutorials
+
+    * `GitHub examples <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/examples>`_
+    * :doc:`Video tutorials <./tutorials/video-tutorials>`
+
+To contribute to the documentation, refer to
+`Contributing to ROCm <https://rocm.docs.amd.com/en/latest/contribute/contributing.html>`_.
+
+You can find licensing information on the
+`Licensing <https://rocm.docs.amd.com/en/latest/about/license.html>`_ page.

projects/rocprofiler-systems/docs/install/install.rst

@@ -0,0 +1,457 @@
+.. meta::
+   :description: ROCm Systems Profiler installation documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, installation, installer, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+*************************************
+ROCm Systems Profiler installation
+*************************************
+
+The following information builds on the guidelines in the :doc:`Quick start <./quick-start>` guide.
+It covers how to install `ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ from
+source or a binary distribution, as well as the :ref:`post-installation-steps`.
+
+If you have problems using ROCm Systems Profiler after installation,
+consult the :ref:`post-installation-troubleshooting` section.
+
+Release links
+========================================
+
+To review and install either the current ROCm Systems Profiler release or earlier releases, use these links:
+
+* Latest ROCm Systems Profiler Release: `<https://github.com/ROCm/rocprofiler-systems/releases/latest>`_
+* All ROCm Systems Profiler Releases: `<https://github.com/ROCm/rocprofiler-systems/releases>`_
+
+Operating system support
+========================================
+
+ROCm Systems Profiler is only supported on Linux. The following distributions are tested in the ROCm Systems Profiler GitHub workflows:
+
+* Ubuntu 20.04
+* Ubuntu 22.04
+* OpenSUSE 15.5
+* OpenSUSE 15.6
+* Red Hat 8.8
+* Red Hat 8.9
+* Red Hat 8.10
+* Red Hat 9.2
+* Red Hat 9.3
+* Red Hat 9.4
+
+Other OS distributions might function but are not supported or tested.
+
+Identifying the operating system
+-----------------------------------
+
+If you are unsure of the operating system and version, the ``/etc/os-release`` and
+``/usr/lib/os-release`` files contain operating system identification data for Linux systems.
+
+.. code-block:: shell
+
+   $ cat /etc/os-release
+
+.. code-block:: shell
+
+   NAME="Ubuntu"
+   VERSION="20.04.4 LTS (Focal Fossa)"
+   ID=ubuntu
+   ...
+   VERSION_ID="20.04"
+   ...
+
+The relevant fields are ``ID`` and the ``VERSION_ID``.
+
+Architecture
+========================================
+
+With regards to instrumentation, at present only AMD64 (x86_64) architectures are tested. However,
+Dyninst supports several more architectures and ROCm Systems Profiler instrumentation may support other
+CPU architectures such as aarch64 and ppc64.
+Other modes of use, such as sampling and causal profiling, are not dependent on Dyninst and therefore
+might be more portable.
+
+Installing ROCm Systems Profiler from binary distributions
+==========================================================
+
+Every ROCm Systems Profiler release provides binary installer scripts of the form:
+
+.. code-block:: shell
+
+   rocprof-sys-{VERSION}-{OS_DISTRIB}-{OS_VERSION}[-ROCm-{ROCM_VERSION}[-{EXTRA}]].sh
+
+For example,
+
+.. code-block:: shell
+
+   rocprof-sys-1.0.0-ubuntu-18.04-OMPT-PAPI-Python3.sh
+   rocprof-sys-1.0.0-ubuntu-18.04-ROCm-405000-OMPT-PAPI-Python3.sh
+   ...
+   rocprof-sys-1.0.0-ubuntu-20.04-ROCm-50000-OMPT-PAPI-Python3.sh
+
+Any of the ``EXTRA`` fields with a CMake build option
+(for example, PAPI, as referenced in a following section) or
+with no link requirements (such as OMPT) have
+self-contained support for these packages.
+
+To install ROCm Systems Profiler using a binary installer script, follow these steps:
+
+#. Download the appropriate binary distribution
+
+   .. code-block:: shell
+
+      wget https://github.com/ROCm/rocprofiler-systems/releases/download/v<VERSION>/<SCRIPT>
+
+#. Create the target installation directory
+
+   .. code-block:: shell
+
+      mkdir /opt/rocprofiler-systems
+
+#. Run the installer script
+
+   .. code-block:: shell
+
+      ./rocprofiler-systems-1.0.0-ubuntu-18.04-ROCm-405000-OMPT-PAPI.sh --prefix=/opt/rocprofiler-systems --exclude-subdir
+
+Building ROCm Systems Profiler from source
+==========================================
+
+ROCm Systems Profiler needs a GCC compiler with full support for C++17 and CMake v3.16 or higher.
+The Clang compiler may be used instead of the GCC compiler if `Dyninst <https://github.com/dyninst/dyninst>`_
+is already installed.
+
+Build requirements
+-----------------------------------
+
+* GCC compiler v7+
+
+  * Older GCC compilers may be supported but are not tested
+  * Clang compilers are generally supported for ROCm Systems Profiler but not Dyninst
+
+* `CMake <https://cmake.org/>`_ v3.16+
+
+  .. note::
+
+     * If the installed version of CMake is too old, installing a new version of CMake can be done through several methods
+     * One of the easiest options is to use the python ``pip`` utility, as follows:
+
+     .. code-block:: shell
+
+        pip install --user 'cmake==3.18.4'
+        export PATH=${HOME}/.local/bin:${PATH}
+
+Required third-party packages
+-----------------------------------
+
+* `Dyninst <https://github.com/dyninst/dyninst>`_ for dynamic or static instrumentation.
+  Dyninst uses the following required and optional components.
+
+  * `TBB <https://github.com/oneapi-src/oneTBB>`_ (required)
+  * `Elfutils <https://sourceware.org/elfutils/>`_ (required)
+  * `Libiberty <https://github.com/gcc-mirror/gcc/tree/master/libiberty>`_ (required)
+  * `Boost <https://www.boost.org/>`_ (required)
+  * `OpenMP <https://www.openmp.org/>`_ (optional)
+
+* `libunwind <https://www.nongnu.org/libunwind/>`_ for call-stack sampling
+
+Any of the third-party packages required by Dyninst, along with Dyninst itself, can be built and installed
+during the ROCm Systems Profiler build. The following list indicates the package, the version,
+the application that requires the package (for example, ROCm Systems Profiler requires Dyninst
+while Dyninst requires TBB), and the CMake option to build the package alongside ROCm Systems Profiler:
+
+.. csv-table::
+   :header: "Third-Party Library", "Minimum Version", "Required By", "CMake Option"
+
+   "Dyninst", "12.0", "ROCm Systems Profiler", "``ROCPROFSYS_BUILD_DYNINST`` (default: OFF)"
+   "Libunwind", "", "ROCm Systems Profiler", "``ROCPROFSYS_BUILD_LIBUNWIND`` (default: ON)"
+   "TBB", "2018.6", "Dyninst", "``ROCPROFSYS_BUILD_TBB`` (default: OFF)"
+   "ElfUtils", "0.178", "Dyninst", "``ROCPROFSYS_BUILD_ELFUTILS`` (default: OFF)"
+   "LibIberty",  "", "Dyninst", "``ROCPROFSYS_BUILD_LIBIBERTY`` (default: OFF)"
+   "Boost",  "1.67.0", "Dyninst", "``ROCPROFSYS_BUILD_BOOST`` (default: OFF)"
+   "OpenMP", "4.x", "Dyninst", ""
+
+Optional third-party packages
+-----------------------------------
+
+* `ROCm <https://rocm.docs.amd.com/projects/install-on-linux/en/latest>`_
+
+  * HIP
+  * AMD SMI Lib for GPU monitoring
+  * ROCprofiler SDK for GPU hardware counters and ROCm tracing
+
+* `PAPI <https://icl.utk.edu/papi/>`_
+* MPI
+
+  * ``ROCPROFSYS_USE_MPI`` enables full MPI support
+  * ``ROCPROFSYS_USE_MPI_HEADERS`` enables wrapping of the dynamically-linked MPI C function calls.
+    (By default, if ROCm Systems Profiler cannot find an OpenMPI MPI distribution, it uses a local copy
+    of the OpenMPI ``mpi.h``.)
+
+* Several optional third-party profiling tools supported by Timemory
+  (for example, `Caliper <https://github.com/LLNL/Caliper>`_, `TAU <https://www.cs.uoregon.edu/research/tau/home.php>`_, CrayPAT, and others)
+
+.. csv-table::
+   :header: "Third-Party Library", "CMake Enable Option", "CMake Build Option"
+   :widths: 15, 45, 40
+
+   "PAPI", "``ROCPROFSYS_USE_PAPI`` (default: ON)", "``ROCPROFSYS_BUILD_PAPI`` (default: ON)"
+   "MPI", "``ROCPROFSYS_USE_MPI`` (default: OFF)", ""
+   "MPI (header-only)", "``ROCPROFSYS_USE_MPI_HEADERS`` (default: ON)", ""
+
+Installing Dyninst
+-----------------------------------
+
+The easiest way to install Dyninst is alongside ROCm Systems Profiler, but it can also be installed using Spack.
+
+Building Dyninst alongside ROCm Systems Profiler
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+To install Dyninst alongside ROCm Systems Profiler, configure ROCm Systems Profiler with ``ROCPROFSYS_BUILD_DYNINST=ON``.
+Depending on the version of Ubuntu, the ``apt`` package manager might have current enough
+versions of the Dyninst Boost, TBB, and LibIberty dependencies
+(use ``apt-get install libtbb-dev libiberty-dev libboost-dev``).
+However, it is possible to request Dyninst to build and install
+its dependencies via ``ROCPROFSYS_BUILD_<DEP>=ON``, as follows:
+
+.. code-block:: shell
+
+   git clone https://github.com/ROCm/rocprofiler-systems.git rocprof-sys-source
+   cmake -B rocprof-sys-build -DROCPROFSYS_BUILD_DYNINST=ON -DROCPROFSYS_BUILD_{TBB,ELFUTILS,BOOST,LIBIBERTY}=ON rocprof-sys-source
+
+where ``-DROCPROFSYS_BUILD_{TBB,BOOST,ELFUTILS,LIBIBERTY}=ON`` is expanded by
+the shell to ``-DROCPROFSYS_BUILD_TBB=ON -DROCPROFSYS_BUILD_BOOST=ON ...``
+
+Installing Dyninst via Spack
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+`Spack <https://github.com/spack/spack>`_ is another option to install Dyninst and its dependencies:
+
+.. code-block:: shell
+
+   git clone https://github.com/spack/spack.git
+   source ./spack/share/spack/setup-env.sh
+   spack compiler find
+   spack external find --all --not-buildable
+   spack spec -I --reuse dyninst
+   spack install --reuse dyninst
+   spack load -r dyninst
+
+.. _cmake-options:
+
+Building and installing ROCm Systems Profiler
+---------------------------------------------
+
+ROCm Systems Profiler has CMake configuration options for MPI support (``ROCPROFSYS_USE_MPI`` or
+``ROCPROFSYS_USE_MPI_HEADERS``),
+ROCm tracing and sampling (``ROCPROFSYS_USE_ROCM``), OpenMP-Tools (``ROCPROFSYS_USE_OMPT``),
+hardware counters via PAPI (``ROCPROFSYS_USE_PAPI``), among other features.
+Various additional features can be enabled via the
+``TIMEMORY_USE_*`` `CMake options <https://timemory.readthedocs.io/en/develop/installation.html#cmake-options>`_.
+Any ``ROCPROFSYS_USE_<VAL>`` option which has a corresponding ``TIMEMORY_USE_<VAL>``
+option means that the Timemory support for this feature has been integrated
+into Perfetto support for ROCm Systems Profiler, for example, ``ROCPROFSYS_USE_PAPI=<VAL>`` also configures
+``TIMEMORY_USE_PAPI=<VAL>``. This means the data that Timemory is able to collect via this package
+is passed along to Perfetto and is displayed when the ``.proto`` file is visualized
+in `the Perfetto UI <https://ui.perfetto.dev>`_.
+
+.. code-block:: shell
+
+   git clone https://github.com/ROCm/rocprofiler-systems.git rocprof-sys-source
+   cmake                                                 \
+       -B rocprof-sys-build                              \
+       -D CMAKE_INSTALL_PREFIX=/opt/rocprofiler-systems  \
+       -D ROCPROFSYS_USE_ROCM=ON                         \
+       -D ROCPROFSYS_USE_PYTHON=ON                       \
+       -D ROCPROFSYS_USE_OMPT=ON                         \
+       -D ROCPROFSYS_USE_MPI_HEADERS=ON                  \
+       -D ROCPROFSYS_BUILD_PAPI=ON                       \
+       -D ROCPROFSYS_BUILD_LIBUNWIND=ON                  \
+       -D ROCPROFSYS_BUILD_DYNINST=ON                    \
+       -D ROCPROFSYS_BUILD_TBB=ON                        \
+       -D ROCPROFSYS_BUILD_BOOST=ON                      \
+       -D ROCPROFSYS_BUILD_ELFUTILS=ON                   \
+       -D ROCPROFSYS_BUILD_LIBIBERTY=ON                  \
+       rocprof-sys-source
+   cmake --build rocprof-sys-build --target all --parallel 8
+   cmake --build rocprof-sys-build --target install
+   source /opt/rocprofiler-systems/share/rocprofiler-systems/setup-env.sh
+
+.. _build-script:
+
+Using the build script
+^^^^^^^^^^^^^^^^^^^^^^
+
+This method automates the CMake process with a script that wraps the CMake
+commands and handles build logic, environment variables, and packaging. Run
+``./scripts/build-release.sh`` with your desired options to generate packages.
+
+Use ``./scripts/build-release.sh --help`` for more information.
+
+.. code-block:: shell-session
+
+   ./scripts/build-release.sh --help
+   Options:
+       --core       [+nopython] [+python]                    Core (Use '+nopython' to build w/o python, use '+python' to python build with python)
+       --mpi        [+nopython] [+python]                    MPI (Use '+nopython' to build w/o python, use '+python' to python build with python)
+       --rocm       [+nopython] [+python]                    ROCm (Use '+nopython' to build w/o python, use '+python' to python build with python)
+       --rocm-mpi   [+nopython] [+python]                    ROCm + MPI (Use '+nopython' to build w/o python, use '+python' to python build with python)
+       --mpi-impl   [openmpi|mpich]                          MPI implementation
+
+       --lto                  [on|off]                       Enable LTO (default: off)
+       --strip                [on|off]                       Strip libraries (default: off)
+       --perfetto-tools       [on|off]                       Install perfetto tools (default: on)
+       --static-libgcc        [on|off]                       Build with static libgcc (default: on)
+       --static-libstdcxx     [on|off]                       Build with static libstdc++ (default: on)
+       --hidden-visibility    [on|off]                       Build with hidden visibility (default: on)
+       --max-threads          N                              Max number of threads supported (default: 2048)
+       --parallel             N                              Number of parallel build jobs (default: 12)
+       --generators           [STGZ][DEB][RPM][+others]      CPack generators (default: stgz deb rpm)
+
+.. _mpi-support-rocprof-sys:
+
+MPI support within ROCm Systems Profiler
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+ROCm Systems Profiler can have full (``ROCPROFSYS_USE_MPI=ON``) or partial (``ROCPROFSYS_USE_MPI_HEADERS=ON``) MPI support.
+The only difference between these two modes is whether or not the results collected
+via Timemory and/or Perfetto can be aggregated into a single
+output file during finalization. When full MPI support is enabled, combining the
+Timemory results always occurs, whereas combining the Perfetto
+results is configurable via the ``ROCPROFSYS_PERFETTO_COMBINE_TRACES`` setting.
+
+The primary benefits of partial or full MPI support are the automatic wrapping
+of MPI functions and the ability
+to label output with suffixes which correspond to the ``MPI_COMM_WORLD`` rank ID
+instead of having to use the system process identifier (i.e. ``PID``).
+In general, it's recommended to use partial MPI support with the OpenMPI
+headers as this is the most portable configuration.
+If full MPI support is selected, make sure your target application is built
+against the same MPI distribution as ROCm Systems Profiler.
+For example, do not build ROCm Systems Profiler with MPICH and use it on a target application built against OpenMPI.
+If partial support is selected, the reason the OpenMPI headers are recommended instead of the MPICH headers is
+because the ``MPI_COMM_WORLD`` in OpenMPI is a pointer to ``ompi_communicator_t`` (8 bytes),
+whereas ``MPI_COMM_WORLD`` in MPICH is an ``int`` (4 bytes). Building ROCm Systems Profiler with partial MPI support
+and the MPICH headers and then using
+ROCm Systems Profiler on an application built against OpenMPI causes a segmentation fault.
+This happens because the value of the ``MPI_COMM_WORLD`` is truncated
+during the function wrapping before being passed along to the underlying MPI function.
+
+ROCm Systems Profiler without ROCm
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+To build ROCm Systems Profiler for use on systems without a GPU or the ROCm runtime, disable ROCm
+support using the CMake configuration option ``ROCPROFSYS_USE_ROCM=OFF``. See :ref:`cmake-options`
+for more information.
+
+Alternatively, use the provided build script with the appropriate options. See :ref:`build-script`.
+For example, to build without ROCm support and create a STGZ installer, use the following command:
+
+.. code-block:: shell
+
+   ./scripts/build-release.sh --core +python --generators STGZ
+
+.. _post-installation-steps:
+
+Post-installation steps
+========================================
+
+After installation, you can optionally configure the ROCm Systems Profiler environment.
+You should also test the executables to confirm ROCm Systems Profiler is correctly installed.
+
+Configure the environment
+-----------------------------------
+
+If environment modules are available and preferred, add them using these commands:
+
+.. code-block:: shell
+
+   module use /opt/rocprofiler-systems/share/modulefiles
+   module load rocprofiler-systems/1.0.0
+
+Alternatively, you can directly source the ``setup-env.sh`` script:
+
+.. code-block:: shell
+
+   source /opt/rocprofiler-systems/share/rocprofiler-systems/setup-env.sh
+
+Test the executables
+-----------------------------------
+
+Successful execution of these commands confirms that the installation does not have any
+issues locating the installed libraries:
+
+.. code-block:: shell
+
+   rocprof-sys-instrument --help
+   rocprof-sys-avail --help
+
+.. note::
+
+   If ROCm support is enabled, you might have to add the path to the ROCm libraries to ``LD_LIBRARY_PATH``,
+   for example, ``export LD_LIBRARY_PATH=/opt/rocm/lib:${LD_LIBRARY_PATH}``.
+
+.. _post-installation-troubleshooting:
+
+Post-installation troubleshooting
+========================================
+
+This section explains how to resolve certain issues that might happen when you first use ROCm Systems Profiler.
+
+Issues with RHEL and SELinux
+----------------------------------------------------
+
+RHEL (Red Hat Enterprise Linux) and related distributions of Linux automatically enable a security feature
+named SELinux (Security-Enhanced Linux) that prevents ROCm Systems Profiler from running.
+This issue applies to any Linux distribution with SELinux installed, including RHEL,
+CentOS, Fedora, and Rocky Linux. The problem can happen with any GPU, or even without a GPU.
+
+The problem occurs after you instrument a program and try to
+run ``rocprof-sys-run`` with the instrumented program.
+
+.. code-block:: shell
+
+   g++ hello.cpp -o hello
+   rocprof-sys-instrument -M sampling -o hello.instr -- ./hello
+   rocprof-sys-run -- ./hello.instr
+
+Instead of successfully running the binary with call-stack sampling,
+ROCm Systems Profiler crashes with a segmentation fault.
+
+.. note::
+
+   If you are physically logged in on the system (not using SSH or a remote connection),
+   the operating system might display an SELinux pop-up warning in the notifications.
+
+To workaround this problem, either disable SELinux or configure it to use a more
+permissive setting.
+
+To avoid this problem for the duration of the current session, run this command
+from the shell:
+
+.. code-block:: shell
+
+   sudo setenforce 0
+
+For a permanent workaround, edit the SELinux configuration file using the command
+``sudo vim /etc/sysconfig/selinux`` and change the ``SELINUX`` setting to
+either ``Permissive`` or ``Disabled``.
+
+.. note::
+
+   Permanently changing the SELinux settings can have security implications.
+   Ensure you review your system security settings before making any changes.
+
+Modifying RPATH details
+----------------------------------------------------
+
+If you're experiencing problems loading your application with an instrumented library,
+then you might have to check and modify the RPATH specified in your application.
+See the section on `troubleshooting RPATHs <../how-to/instrumenting-rewriting-binary-application.html#rpath-troubleshooting>`_
+for further details.
+
+Configuring PAPI to collect hardware counters
+----------------------------------------------------
+
+To use PAPI to collect the majority of hardware counters, ensure
+the ``/proc/sys/kernel/perf_event_paranoid`` setting has a value less than or equal to ``2``.
+For more information, see the :ref:`rocprof-sys_papi_events` section.
+

projects/rocprofiler-systems/docs/install/quick-start.rst

@@ -0,0 +1,51 @@
+.. meta::
+   :description: ROCm Systems Profiler quick start documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, quick start, getting started, quick install, tracking, visualization, tool, Instinct, accelerator, AMD
+
+*************************************
+ROCm Systems Profiler quick start
+*************************************
+
+To install ROCm Systems Profiler, download the
+`ROCm Systems Profiler installer <https://github.com/ROCm/rocprofiler-systems/releases/latest/download/rocprofiler-systems-install.py>`_
+and specify ``--prefix <install-directory>``. The script attempts to auto-detect
+the appropriate OS distribution and version. To include AMD ROCm Software support,
+specify ``--rocm X.Y``, where ``X`` is the ROCm major
+version and ``Y`` is the ROCm minor version, for example, ``--rocm 6.3``.
+
+.. code-block:: shell
+
+   wget https://github.com/ROCm/rocprofiler-systems/releases/latest/download/rocprofiler-systems-install.py
+   python3 ./rocprofiler-systems-install.py --prefix /opt/rocprofiler-systems --rocm 6.3
+
+This script supports installation on Ubuntu, OpenSUSE, Red Hat, Debian, CentOS, and Fedora.
+If the target OS is compatible with one of the operating system versions listed in
+the comprehensive :doc:`Installation guidelines <./install>`,
+specify ``-d <DISTRO> -v <VERSION>``. For example, if the OS is compatible with Ubuntu 22.04, pass
+``-d ubuntu -v 22.04`` to the script.
+
+Install via package manager
+============================
+
+   If you have ROCm version 6.3 or higher installed, you can use the
+   package manager to install a pre-built copy of ROCm Systems Profiler.
+
+.. tab-set::
+
+   .. tab-item:: Ubuntu
+
+      .. code-block:: shell
+
+         $ sudo apt install rocprofiler-systems
+
+   .. tab-item:: Red Hat Enterprise Linux
+
+      .. code-block:: shell
+
+         $ sudo dnf install rocprofiler-systems
+
+   .. tab-item:: SUSE Linux Enterprise Server
+
+      .. code-block:: shell
+
+         $ sudo zypper install rocprofiler-systems

projects/rocprofiler-systems/docs/license.rst

@@ -0,0 +1,8 @@
+.. meta::
+   :description: ROCm Systems Profiler license
+
+*******
+License
+*******
+
+.. include:: ../LICENSE

projects/rocprofiler-systems/docs/reference/development-guide.rst

@@ -0,0 +1,412 @@
+.. meta::
+   :description: ROCm Systems Profiler development documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, development, developers guide, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Development guide
+****************************************************
+
+This guide discusses the `ROCm Systems Profiler <https://github.com/ROCm/rocprofiler-systems>`_ design.
+It includes a list of the executables and libraries, along with a discussion of the application's
+memory, sampling, and time-window constraint models.
+
+Executables
+========================================
+
+This section lists the ROCm Systems Profiler executables.
+
+rocprof-sys-avail: `source/bin/rocprof-sys-avail <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/bin/rocprof-sys-avail>`_
+-----------------------------------------------------------------------------------------------------------------------------------------------
+
+The ``main`` routine of ``rocprof-sys-avail`` has three important sections:
+
+* Printing components
+* Printing options
+* Printing hardware counters
+
+rocprof-sys-sample: `source/bin/rocprof-sys-sample <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/bin/rocprof-sys-sample>`_
+--------------------------------------------------------------------------------------------------------------------------------------------------
+
+* Requires a command-line format of ``rocprof-sys-sample <options> -- <command> <command-args>``
+* Translates command-line options into environment variables
+* Adds ``librocprof-sys-dl.so`` to ``LD_PRELOAD``
+* Is launched by using ``execvpe`` with ``<command> <command-args>`` and a modified environment
+
+rocprof-sys-casual: `source/bin/rocprof-sys-causal <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/bin/rocprof-sys-causal>`_
+---------------------------------------------------------------------------------------------------------------------------------------------------
+
+When there is exactly one causal profiling configuration variant (which enables debugging),
+``rocprof-sys-casual`` has a nearly identical design to ``rocprof-sys-sample``
+
+When the command-line options produce more than one causal profiling configuration variant,
+the following actions take place for each variant:
+
+* ``rocprof-sys-causal`` calls ``fork()``
+* the child process launches ``<command> <command-args>`` using ``execvpe``, which modifies the environment for the variant
+* the parent process waits for the child process to finish
+
+rocprof-sys-instrument: `source/bin/rocprof-sys-instrument <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/bin/rocprof-sys-instrument>`_
+--------------------------------------------------------------------------------------------------------------------------------------------------------------
+
+* Requires a command-line format of ``rocprof-sys-instrument <options> -- <command> <command-args>``
+* Allows the user to provide options specifying whether to perform runtime instrumentation, use binary rewrite, or
+  attach to process
+* Either opens the instrumentation target (for binary rewrite), launches the target and stops it
+  before it starts executing ``main``, or attaches to a running executable and pauses it
+* Finds all functions in the targets
+* Finds ``librocprof-sys-dl`` and locates the functions
+* Iterates over and instruments all the functions, provided they satisfy the
+  defined criteria (such as a minimum number of instructions)
+
+  * See the ``module_function`` class
+
+* Until this point, the workflow has been the same for the different options,
+  but it diverges after instrumentation is complete:
+
+  * For a binary rewrite: it produces a new instrumented binary and exits
+  * For runtime instrumentation or attaching to a process: it instructs the application
+    to resume and then waits for it to exit
+
+Libraries
+========================================
+
+Common library: `source/lib/common <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/lib/common>`_
+--------------------------------------------------------------------------------------------------------------------------------
+
+* General header-only functionality used in multiple executables and/or libraries.
+* Not installed or exported outside of the build tree.
+
+Core library: `source/lib/core <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/lib/core>`_
+--------------------------------------------------------------------------------------------------------------------------------
+
+* Static PIC library with functionality that does not depend on any components.
+* Not installed or exported outside of the build tree.
+
+Binary library: `source/lib/binary <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/lib/binary>`_
+--------------------------------------------------------------------------------------------------------------------------------
+
+* Static PIC library with functionality for reading/analyzing binary info.
+* Mostly used by the causal profiling sections of ``librocprof-sys``.
+* Not installed or exported outside of the build tree.
+
+librocprof-sys: `source/lib/rocprof-sys <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/lib/rocprof-sys>`_
+--------------------------------------------------------------------------------------------------------------------------------
+
+This is the main library encapsulating all the capabilities.
+
+librocprof-sys-dl: `source/lib/rocprof-sys-dl <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/lib/rocprof-sys-dl>`_
+-----------------------------------------------------------------------------------------------------------------------------------------
+
+This is a lightweight, front-end library for ``librocprof-sys`` which serves three primary purposes:
+
+* Dramatically speeds up instrumentation time compared to using ``librocprof-sys`` directly because
+  Dyninst must parse the entire library in order to find the instrumentation functions
+  (a ``dlopen`` call is made on ``librocprof-sys`` when the instrumentation functions get called)
+* Prevents re-entry if ``librocprof-sys`` calls an instrumented function internally
+* Coordinates communication between ``librocprof-sys-user`` and ``librocprof-sys``
+
+librocprof-sys-user: `source/lib/rocprof-sys-user <https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/source/lib/rocprof-sys-user>`_
+-----------------------------------------------------------------------------------------------------------------------------------------------
+
+* Provides a set of functions and types for the users to add to their code, for example,
+  disabling data collection globally or on a specific thread or
+  user-defined region
+* If ``librocprof-sys-dl`` is not loaded, the user API is effectively a set of no-op function calls.
+
+Testing tools
+========================================
+
+* `CDash Testing Dashboard <https://my.cdash.org/index.php?project=rocprofiler-systems>`_ (requires a login)
+
+Components
+========================================
+
+Most measurements and capabilities are encapsulated into a "component" with the following definitions:
+
+Measurement
+   A recording of some data relevant to performance, for instance, the current call-stack,
+   hardware counter values, current memory usage, or timestamp
+
+Capability
+   Handles the implementation or orchestration of some feature which is used
+   to collect measurements, for example, a component which handles setting up function wrappers
+   around various functions such as ``pthread_create`` or ``MPI_Init``.
+
+Components are designed to either hold no data at all or only the data for both an instantaneous
+measurement and a phase measurement.
+
+Components which store data typically implement a static ``record()`` function
+for getting a record of the measurement,
+``start()`` and ``stop()`` member functions for calculating a phase measurement,
+and a ``sample()`` member function for storing an
+instantaneous measurement. In reality, there are several more "standard" functions
+but these are the most commonly-used ones.
+
+Components which do not store data might also have ``start()``, ``stop()``, and ``sample()``
+functions. However, components which
+implement function wrappers typically provide a call operator or ``audit(...)``
+functions. These are invoked with the
+wrapped function's arguments before the wrapped function gets called and with the return value
+after the wrapped function gets called.
+
+.. note::
+
+   The goal of this design is to provide relatively small and resuable lightweight objects
+   for recording measurements and implementing capabilities.
+
+Wall-clock component example
+--------------------------------------
+
+A component for computing the elapsed wall-clock time looks like this:
+
+.. code-block:: cpp
+
+   struct wall_clock
+   {
+      using value_type = int64_t;
+
+      static value_type record() noexcept
+      {
+         return std::chrono::steady_clock::now().time_since_epoch().count();
+      }
+
+      void sample() noexcept
+      {
+         value = record();
+      }
+
+      void start() noexcept
+      {
+         value = record();
+      }
+
+      void stop() noexcept
+      {
+         auto _start_value = value;
+         value = record();
+         accum += (value - _start_value);
+      }
+
+   private:
+      int64_t value = 0;
+      int64_t accum = 0;
+   };
+
+Function wrapper component example
+--------------------------------------
+
+A component which implements wrappers around ``fork()`` and ``exit(int)`` (and stores no data)
+could look like this:
+
+.. code-block:: cpp
+
+   struct function_wrapper
+   {
+      pid_t operator()(const gotcha_data&, pid_t (*real_fork)())
+      {
+         // disable all collection before forking
+         categories::disable_categories(config::get_enabled_categories());
+
+         auto _pid_v = real_fork();
+
+         // only re-enable collection on parent process
+         if(_pid_v != 0)
+               categories::enable_categories(config::get_enabled_categories());
+
+         return _pid_v;
+      }
+
+      void operator()(const gotcha_data&, void (*real_exit)(int), int _exit_code)
+      {
+         // catch the call to exit and finalize before truly exiting
+         rocprofsys_finalize();
+
+         real_exit(_exit_code);
+      }
+   };
+
+Component member functions
+--------------------------------------
+
+There are no real restrictions or requirements on the member functions a component needs to provide.
+Unless the component is being used directly, the invocation of component member functions via a "component bundler"
+(provided by Timemory) makes extensive use of template metaprogramming concepts. This finds the best match, if any,
+for calling a component's member function. This is a bit easier to demonstrate using an example:
+
+.. code-block:: cpp
+
+   struct foo
+   {
+      void sample() { puts("foo::sample()"); }
+   };
+
+   struct bar
+   {
+      void sample(int) { puts("bar::sample(int)"); }
+   };
+
+   struct spam
+   {
+      void start(int) { puts("spam::start()"); }
+      void stop()     { puts("spam::stop()"); }
+   };
+
+   int main()
+   {
+      auto _bundle = component_tuple<foo, bar, spam>{ "main" };
+
+      puts("A");
+      _bundle.start();
+
+      puts("B");
+      _bundle.sample(10);
+
+      puts("C");
+      _bundle.sample();
+
+      puts("D");
+      _bundle.stop();
+   }
+
+When the preceding code runs, the following messages are printed:
+
+.. code-block:: shell
+
+   A
+   spam::start()
+   B
+   foo::sample()
+   bar::sample(int)
+   C
+   foo::sample()
+   D
+   spam::stop()
+
+In section A, the bundle determined that only the ``spam`` object has a ``start`` function. Since this is determined
+via template metaprogramming instead of dynamic polymorphism, this effectively omits any code related to
+the ``foo`` or ``bar`` objects. In section B, because the integer ``10`` is passed to the bundle,
+the bundle forwards this value to ``bar::sample(int)`` after it invokes ``foo::sample()``. ``foo::sample()`` is
+invoked because the bundle recognizes that the call to the ``sample`` member function is still possible without
+the argument.
+
+Memory model
+========================================
+
+Collected data is generally handled in one of the three following ways:
+
+* It is handed directly to, and stored by, Perfetto
+* It is managed implicitly by Timemory and accessed as needed
+* As thread-local data
+
+In general, only instrumentation for relatively simple data is directly passed to
+Perfetto and/or Timemory during runtime.
+For example, the callbacks from binary instrumentation, user API instrumentation,
+and rocprofiler-sdk directly invoke
+calls to Perfetto or Timemory's storage model. Otherwise, the data is stored
+by ROCm Systems Profiler in the thread-data model
+which is more persistent than simply using ``thread_local`` static data, which gets deleted
+when the thread stops.
+
+Thread identification
+--------------------------------------
+
+Each CPU thread is assigned two integral identifiers. One identifier, the ``internal_value``, is
+atomically incremented every time a new thread is created.
+The other identifier, known as the ``sequent_value``, tries to account for the fact that ROCm Systems Profiler, Perfetto, ROCm, and other applications
+start background threads. When a thread is created as a by-product of ROCm Systems Profiler,
+the index is offset by a large value. This serves
+two purposes:
+
+* Accessing the data for threads created by the user is closer in memory
+* When log messages are printed, the index approximately correlates to the order of thread creation from the user's perspective.
+
+The ``sequent_value`` identifier is typically used to access the thread-data.
+
+Thread-data class
+--------------------------------------
+
+Currently, most thread data is effectively stored in a static
+``std::array<std::unique_ptr<T>, ROCPROFSYS_MAX_THREADS>`` instance.
+``ROCPROFSYS_MAX_THREADS`` is a value defined a compile-time and set to ``2048``
+for release builds. During finalization,
+ROCm Systems Profiler iterates through the thread-data and transforms that data
+into something that can be passed along to Perfetto and/or Timemory.
+The downside of the current model is that if the user exceeds ``ROCPROFSYS_MAX_THREADS``,
+a segmentation fault occurs. To fix this issue,
+a new model is being adopted which has all the benefits of this model
+but permits dynamic expansion.
+
+Sampling model
+========================================
+
+The general structure for the sampling is within Timemory (``source/timemory/sampling``).
+Currently, all sampling is done per-thread
+via POSIX timers. ROCm Systems Profiler supports both a real-time timer and a CPU-time timer.
+Both have adjustable frequencies, delays, and durations.
+By default, only CPU-time sampling is enabled. Initial settings are inherited from
+the settings starting with ``ROCPROFSYS_SAMPLING_``.
+
+For each type of timer, timer-specific settings can be used to
+override the common and inherited timer settings.
+These settings begin with ``ROCPROFSYS_SAMPLING_CPUTIME`` for the CPU-time sampler
+and ``ROCPROFSYS_SAMPLING_REALTIME`` for
+the real-time sampler. For example, ``ROCPROFSYS_SAMPLING_FREQ=500`` initially sets the
+sampling frequency to 500 interrupts per second. Adding the setting ``ROCPROFSYS_SAMPLING_REALTIME_FREQ=10``
+lowers the sampling frequency for the real-time sampler
+to 10 interrupts per second of real-time.
+
+The ROCm Systems Profiler-specific implementation can be found in
+`source/lib/rocprof-sys/library/sampling.cpp <https://github.com/ROCm/rocprofiler-systems/blob/amd-mainline/source/lib/rocprof-sys/library/sampling.cpp>`_.
+Within `sampling.cpp <https://github.com/ROCm/rocprofiler-systems/blob/amd-mainline/source/lib/rocprof-sys/library/sampling.cpp>`_,
+there is a bundle of three sampling components:
+
+* `backtrace_timestamp <https://github.com/ROCm/rocprofiler-systems/blob/amd-mainline/source/lib/rocprof-sys/library/components/backtrace_timestamp.hpp>`_ simply
+  records the wall-clock time of the sample.
+* `backtrace <https://github.com/ROCm/rocprofiler-systems/blob/amd-mainline/source/lib/rocprof-sys/library/components/backtrace.hpp>`_
+  records the call-stack via libunwind.
+* `backtrace_metrics <https://github.com/ROCm/rocprofiler-systems/blob/amd-mainline/source/lib/rocprof-sys/library/components/backtrace_metrics.hpp>`_
+  records the sample metrics, such as peak RSS and the hardware counters.
+
+These three components are bundled together in
+a tuple-like ``struct`` (``tuple<backtrace_timestamp, backtrace, backtrace_metrics>``).
+A buffer of at least 1024 instances of this tuple is mapped using ``mmap``
+per-thread. When this buffer is full,
+the sampler hands the buffer off to its allocator thread and maps a new buffer with ``mmap``
+before taking the next sample. The allocator thread takes this data
+and either dynamically stores it in memory or writes it to a file depending on the
+value of ``ROCPROFSYS_USE_TEMPORARY_FILES``.
+This schema avoids all allocations in the signal handler, lets the data grow
+dynamically, avoids potentially slow I/O within the signal handler, and also enables
+the capability of avoiding I/O altogether.
+The maximum number of samplers handled by each allocator is governed by the
+``ROCPROFSYS_SAMPLING_ALLOCATOR_SIZE`` setting (the default is eight). Whenever an allocator
+has reached its limit,
+a new internal thread is created to handle the new samplers.
+
+Time-window constraint model
+========================================
+
+With the recent introduction of tracing delay and duration, the
+`constraint namespace <https://github.com/ROCm/rocprofiler-systems/blob/amd-mainline/source/lib/core/constraint.hpp>`_
+was introduced to improve the management of delays and duration limits for
+data collection. The ``spec`` class accepts a clock identifier, a delay value, a duration value, and an
+integer indicating how many times to repeat the delay and duration cycle. It is therefore
+possible to perform tasks such as periodically enabling tracing for brief periods
+of time in between long periods without data collection while the application runs. The
+syntax follows the format ``clock_identifier:delay:capture_duration:cycles``, so a value of
+``10:1:3`` for the last three parameters represents the following sequence of operations:
+
+* Ten seconds where no data is collected, then one second where it is
+* Ten seconds where no data is collected, then one second where it is
+* Ten seconds where no data is collected, then one second where it is
+* Stop
+
+As another example, ``ROCPROFSYS_TRACE_PERIODS = realtime:10:1:5 process_cputime:10:2:20`` translates
+to this sequence:
+
+* Five cycles of: no data collection for ten seconds of real-time followed by one second of data collection
+* Twenty cycles of: no data collection for ten seconds of process CPU time followed by two CPU-time seconds of data collection
+
+Eventually, the goal is to migrate all subsets of data collection which currently support
+more rudimentary models of time window constraints, such as process sampling and causal profiling,
+to this model.

projects/rocprofiler-systems/docs/reference/rocprof-sys-glossary.rst

@@ -0,0 +1,102 @@
+.. meta::
+   :description: ROCm Systems Profiler glossary and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, glossary, terminology, profiler, tracking, visualization, tool, Instinct, accelerator, AMD
+
+********
+Glossary
+********
+
+This topic explains the terminology necessary to use ROCm Systems Profiler.
+The list below provides a basic glossary for those who
+are new to binary instrumentation. It also clarifies ambiguities
+when certain terms have different
+contextual meanings, for example, the ROCm Systems Profiler meaning of the term "module"
+when instrumenting Python.
+
+Binary
+  A file written in the Executable and Linkable Format (ELF). This is the standard file
+  format for executable files, shared libraries, etc.
+
+Binary instrumentation
+  Inserting callbacks to instrumentation into an existing binary. This can be performed
+  statically or dynamically.
+
+Static binary instrumentation
+  Loads an existing binary, determines instrumentation points, and generates a new binary
+  with instrumentation directly embedded. It is applicable to executables and libraries but
+  limited to only the functions defined in the binary. This is also known as **Binary rewrite**.
+
+Dynamic binary instrumentation
+  Loads an existing binary into memory, inserts instrumentation, and runs the binary.
+  It is limited to executables but is capable of instrumenting linked libraries.
+  This is also known as **Runtime instrumentation**.
+
+Statistical sampling
+  At periodic intervals, the application is paused and the current call-stack of the CPU
+  is recorded along with various other metrics. It uses timers that measure either
+  (A) real clock time or (B) the CPU time used by the current thread and the CPU time
+  expended on behalf of the thread by the system. This is also known as simply **sampling**.
+
+  Sampling rate
+    * The period at which (A) or (B) are triggered (in units of ``# interrupts / second``)
+    * Higher values increase the number of samples
+
+  Sampling delay
+    * How long to wait before (A) and (B) begin triggering at their designated rate
+
+  Sampling duration
+    * The amount of time (in real-time) after the start of the application to record samples.
+    * After this time limit has been reached, no more samples are recorded.
+
+Process sampling
+  At periodic (real-time) intervals, a background thread records global metrics without
+  interrupting the current process. These metrics include, but are not limited to:
+  CPU frequency, CPU memory high-water mark (i.e. peak memory usage), GPU temperature,
+  and GPU power usage.
+
+  Sampling rate
+    * The real-time period for recording metrics (in units of ``# measurements / second``)
+    * Higher values increase the number of samples
+
+  Sampling delay
+    * How long to wait (in real-time) before recording samples
+
+  Sampling duration
+    * The amount of time (in real-time) after the start of the application to record samples.
+    * After this time limit has been reached, no more samples are recorded.
+
+Module
+  With respect to binary instrumentation, a module is defined as either the filename
+  (such as ``foo.c``) or library name (``libfoo.so``) which contains the definition
+  of one or more functions.
+
+  With respect to Python instrumentation, a module is defined as the **file** which contains
+  the definition of one or more functions. The full path to this file typically contains the
+  name of the "Python module".
+
+Basic block
+  A straight-line code sequence with no branches in (except for the entry) and
+  no branches out (except for the exit).
+
+Address range
+  The instructions for a function in a binary start at certain address with the ELF file
+  and end at a certain address. The range is ``end - start``.
+
+  The address range is a decent approximation for the "cost" of a function.
+  For example, a larger address range approximately equates to more instructions.
+
+Instrumentation traps
+  On the x86 architecture, because instructions are of variable size, an instruction
+  might be too small for Dyninst to replace it with the normal code sequence
+  used to call instrumentation. When instrumentation is placed at points other
+  than subroutine entry, exit, or call points, traps may be used to ensure
+  the instrumentation fits. (By default, ``rocprof-sys-instrument`` avoids instrumentation
+  which requires a trap.)
+
+Overlapping functions
+  Due to language constructs or compiler optimizations, it might be possible for
+  multiple functions to overlap (that is, share part of the same function body)
+  or for a single function to have multiple entry points. In practice, it's
+  impossible to determine the difference between multiple overlapping functions
+  and a single function with multiple entry points. (By default, ``rocprof-sys-instrument``
+  avoids instrumenting overlapping functions.)

projects/rocprofiler-systems/docs/sphinx/_toc.yml.in

@@ -0,0 +1,78 @@
+# Anywhere {branch} is used, the branch name will be substituted.
+# These comments will also be removed.
+defaults:
+  numbered: False
+  maxdepth: 6
+root: index
+subtrees:
+  - entries:
+    - file: what-is-rocprof-sys.rst
+
+  - caption: Install
+    entries:
+    - file: install/quick-start.rst
+      title: ROCm Systems Profiler quick start
+    - file: install/install.rst
+      title: ROCm Systems Profiler installation guide
+
+  - caption: How to
+    entries:
+    - file: how-to/configuring-validating-environment.rst
+      title: Configuring the environment
+      subtrees:
+      - entries:
+        - file: how-to/configuring-runtime-options.rst
+          title: Configuring runtime options
+    - file: how-to/general-tips-using-rocprof-sys.rst
+      title: Profiling
+      subtrees:
+      - entries:
+        - file: how-to/sampling-call-stack.rst
+          title: Sampling the call stack
+        - file: how-to/instrumenting-rewriting-binary-application.rst
+          title: Instrumenting and rewriting a binary application
+        - file: how-to/performing-causal-profiling.rst
+          title: Performing causal profiling
+        - file: how-to/profiling-python-scripts.rst
+          title: Profiling Python scripts
+        - file: how-to/nic-profiling.rst
+          title: Network performance profiling
+        - file: how-to/vcn-jpeg-sampling.rst
+          title: VCN and JPEG sampling and tracing
+    - file: how-to/understanding-rocprof-sys-output.rst
+      title: Understanding the output
+    - file: how-to/using-rocprof-sys-api.rst
+      title: Using the ROCm Systems Profiler API
+
+  - caption: Conceptual
+    entries:
+    - file: conceptual/rocprof-sys-feature-set.rst
+      title: Features and use cases
+    - file: conceptual/data-collection-modes.rst
+      title: Data collection modes
+
+  - caption: Reference
+    entries:
+    - file: reference/development-guide.rst
+      title: Development guide
+    - file: reference/rocprof-sys-glossary.rst
+      title: Glossary
+    - file: doxygen/html/files
+      title: API library
+    - file: doxygen/html/functions
+      title: Class member functions
+    - file: doxygen/html/globals
+      title: Globals
+    - file: doxygen/html/annotated
+      title: Classes, structures, and interfaces
+
+  - caption: Tutorials
+    entries:
+    - url: https://github.com/ROCm/rocprofiler-systems/tree/amd-mainline/examples
+      title: GitHub examples
+    - file: tutorials/video-tutorials.rst
+      title: Video tutorials
+
+  - caption: About
+    entries:
+    - file: license.rst

projects/rocprofiler-systems/docs/sphinx/requirements.in

@@ -0,0 +1 @@
+rocm-docs-core[api_reference]==1.21.1

projects/rocprofiler-systems/docs/sphinx/requirements.txt

@@ -0,0 +1,318 @@
+#
+# This file is autogenerated by pip-compile with Python 3.10
+# by the following command:
+#
+#    pip-compile requirements.in
+#
+accessible-pygments==0.0.5
+    # via pydata-sphinx-theme
+alabaster==1.0.0
+    # via sphinx
+asttokens==3.0.0
+    # via stack-data
+attrs==24.3.0
+    # via
+    #   jsonschema
+    #   jupyter-cache
+    #   referencing
+babel==2.16.0
+    # via
+    #   pydata-sphinx-theme
+    #   sphinx
+beautifulsoup4==4.12.3
+    # via pydata-sphinx-theme
+breathe==4.35.0
+    # via rocm-docs-core
+certifi==2024.8.30
+    # via requests
+cffi==1.17.1
+    # via
+    #   cryptography
+    #   pynacl
+charset-normalizer==3.3.2
+    # via requests
+click==8.1.7
+    # via
+    #   click-log
+    #   doxysphinx
+    #   jupyter-cache
+    #   sphinx-external-toc
+click-log==0.4.0
+    # via doxysphinx
+comm==0.2.2
+    # via ipykernel
+contourpy==1.3.1
+    # via matplotlib
+cryptography==44.0.1
+    # via pyjwt
+cycler==0.12.1
+    # via matplotlib
+debugpy==1.8.12
+    # via ipykernel
+decorator==5.1.1
+    # via ipython
+deprecated==1.2.14
+    # via pygithub
+docutils==0.21.2
+    # via
+    #   breathe
+    #   myst-parser
+    #   pydata-sphinx-theme
+    #   sphinx
+doxysphinx==3.3.12
+    # via rocm-docs-core
+exceptiongroup==1.2.2
+    # via ipython
+executing==2.1.0
+    # via stack-data
+fastjsonschema==2.20.0
+    # via
+    #   nbformat
+    #   rocm-docs-core
+fonttools==4.55.0
+    # via matplotlib
+gitdb==4.0.11
+    # via gitpython
+gitpython==3.1.43
+    # via rocm-docs-core
+greenlet==3.1.1
+    # via sqlalchemy
+idna==3.10
+    # via requests
+imagesize==1.4.1
+    # via sphinx
+importlib-metadata==8.6.1
+    # via
+    #   jupyter-cache
+    #   myst-nb
+ipykernel==6.29.5
+    # via myst-nb
+ipython==8.31.0
+    # via
+    #   ipykernel
+    #   myst-nb
+jedi==0.19.2
+    # via ipython
+jinja2==3.1.6
+    # via
+    #   myst-parser
+    #   sphinx
+jsonschema==4.23.0
+    # via nbformat
+jsonschema-specifications==2024.10.1
+    # via jsonschema
+jupyter-cache==1.0.1
+    # via myst-nb
+jupyter-client==8.6.3
+    # via
+    #   ipykernel
+    #   nbclient
+jupyter-core==5.7.2
+    # via
+    #   ipykernel
+    #   jupyter-client
+    #   nbclient
+    #   nbformat
+kiwisolver==1.4.7
+    # via matplotlib
+libsass==0.22.0
+    # via doxysphinx
+lxml==5.2.1
+    # via doxysphinx
+markdown-it-py==3.0.0
+    # via
+    #   mdit-py-plugins
+    #   myst-parser
+markupsafe==2.1.5
+    # via jinja2
+matplotlib==3.9.2
+    # via doxysphinx
+matplotlib-inline==0.1.7
+    # via
+    #   ipykernel
+    #   ipython
+mdit-py-plugins==0.4.2
+    # via myst-parser
+mdurl==0.1.2
+    # via markdown-it-py
+mpire==2.10.2
+    # via doxysphinx
+myst-nb==1.1.2
+    # via rocm-docs-core
+myst-parser==4.0.0
+    # via myst-nb
+nbclient==0.10.2
+    # via
+    #   jupyter-cache
+    #   myst-nb
+nbformat==5.10.4
+    # via
+    #   jupyter-cache
+    #   myst-nb
+    #   nbclient
+nest-asyncio==1.6.0
+    # via ipykernel
+numpy==1.26.4
+    # via
+    #   contourpy
+    #   doxysphinx
+    #   matplotlib
+packaging==24.1
+    # via
+    #   ipykernel
+    #   matplotlib
+    #   sphinx
+parso==0.8.4
+    # via jedi
+pexpect==4.9.0
+    # via ipython
+pillow==11.0.0
+    # via matplotlib
+platformdirs==4.3.6
+    # via jupyter-core
+prompt-toolkit==3.0.50
+    # via ipython
+psutil==6.1.1
+    # via ipykernel
+ptyprocess==0.7.0
+    # via pexpect
+pure-eval==0.2.3
+    # via stack-data
+pycparser==2.22
+    # via cffi
+pydata-sphinx-theme==0.16.1
+    # via
+    #   rocm-docs-core
+    #   sphinx-book-theme
+pygithub==2.4.0
+    # via rocm-docs-core
+pygments==2.18.0
+    # via
+    #   accessible-pygments
+    #   ipython
+    #   mpire
+    #   pydata-sphinx-theme
+    #   sphinx
+pyjson5==1.6.7
+    # via doxysphinx
+pyjwt[crypto]==2.9.0
+    # via pygithub
+pynacl==1.5.0
+    # via pygithub
+pyparsing==3.2.0
+    # via
+    #   doxysphinx
+    #   matplotlib
+python-dateutil==2.9.0.post0
+    # via
+    #   jupyter-client
+    #   matplotlib
+pyyaml==6.0.2
+    # via
+    #   jupyter-cache
+    #   myst-nb
+    #   myst-parser
+    #   rocm-docs-core
+    #   sphinx-external-toc
+pyzmq==26.2.0
+    # via
+    #   ipykernel
+    #   jupyter-client
+referencing==0.36.1
+    # via
+    #   jsonschema
+    #   jsonschema-specifications
+requests==2.32.3
+    # via
+    #   pygithub
+    #   sphinx
+rocm-docs-core[api-reference]==1.21.1
+    # via -r requirements.in
+rpds-py==0.22.3
+    # via
+    #   jsonschema
+    #   referencing
+six==1.16.0
+    # via python-dateutil
+smmap==5.0.1
+    # via gitdb
+snowballstemmer==2.2.0
+    # via sphinx
+soupsieve==2.6
+    # via beautifulsoup4
+sphinx==8.0.2
+    # via
+    #   breathe
+    #   myst-nb
+    #   myst-parser
+    #   pydata-sphinx-theme
+    #   rocm-docs-core
+    #   sphinx-book-theme
+    #   sphinx-copybutton
+    #   sphinx-design
+    #   sphinx-external-toc
+    #   sphinx-notfound-page
+sphinx-book-theme==1.1.3
+    # via rocm-docs-core
+sphinx-copybutton==0.5.2
+    # via rocm-docs-core
+sphinx-design==0.6.1
+    # via rocm-docs-core
+sphinx-external-toc==1.0.1
+    # via rocm-docs-core
+sphinx-notfound-page==1.0.4
+    # via rocm-docs-core
+sphinxcontrib-applehelp==2.0.0
+    # via sphinx
+sphinxcontrib-devhelp==2.0.0
+    # via sphinx
+sphinxcontrib-htmlhelp==2.1.0
+    # via sphinx
+sphinxcontrib-jsmath==1.0.1
+    # via sphinx
+sphinxcontrib-qthelp==2.0.0
+    # via sphinx
+sphinxcontrib-serializinghtml==2.0.0
+    # via sphinx
+sqlalchemy==2.0.37
+    # via jupyter-cache
+stack-data==0.6.3
+    # via ipython
+tabulate==0.9.0
+    # via jupyter-cache
+tomli==2.0.2
+    # via sphinx
+tornado==6.5.1
+    # via
+    #   ipykernel
+    #   jupyter-client
+tqdm==4.67.1
+    # via mpire
+traitlets==5.14.3
+    # via
+    #   comm
+    #   ipykernel
+    #   ipython
+    #   jupyter-client
+    #   jupyter-core
+    #   matplotlib-inline
+    #   nbclient
+    #   nbformat
+typing-extensions==4.12.2
+    # via
+    #   ipython
+    #   myst-nb
+    #   pydata-sphinx-theme
+    #   pygithub
+    #   referencing
+    #   sqlalchemy
+urllib3==2.2.3
+    # via
+    #   pygithub
+    #   requests
+wcwidth==0.2.13
+    # via prompt-toolkit
+wrapt==1.16.0
+    # via deprecated
+zipp==3.21.0
+    # via importlib-metadata

projects/rocprofiler-systems/docs/tutorials/video-tutorials.rst

@@ -0,0 +1,38 @@
+.. meta::
+   :description: ROCm Systems Profiler video documentation and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, video, tutorial, demonstration, tracking, visualization, tool, Instinct, accelerator, AMD
+
+****************************************************
+Video tutorials
+****************************************************
+
+The following video tutorials provide a visual guide to using ROCm Systems Profiler.
+They were recorded using the former name of the tool, Omnitrace, but the content is still applicable.
+
+Installing a binary release
+========================================
+
+.. raw:: html
+
+   <p align="center"><iframe width="560" height="315" src="https://www.youtube.com/embed/gKtNCKf1IXA?modestbranding=1" title="YouTube video player" frameborder="0" allow="accelerometer; clipboard-write; encrypted-media; gyroscope; picture-in-picture" allowfullscreen></iframe></p>
+
+Instrumenting a binary
+========================================
+
+.. raw:: html
+
+   <p align="center"><iframe width="560" height="315" src="https://www.youtube.com/embed/2B0gRr3FygQ?modestbranding=1" title="YouTube video player" frameborder="0" allow="accelerometer; clipboard-write; encrypted-media; gyroscope; picture-in-picture" allowfullscreen></iframe></p>
+
+Writing a ROCm Systems Profiler configuration file
+==================================================
+
+.. raw:: html
+
+   <p align="center"><iframe width="560" height="315" src="https://www.youtube.com/embed/oG_fPYx9_gs?modestbranding=1" title="YouTube video player" frameborder="0" allow="accelerometer; clipboard-write; encrypted-media; gyroscope; picture-in-picture" allowfullscreen></iframe></p>
+
+Visualization and features of Perfetto traces
+=============================================
+
+.. raw:: html
+
+   <p align="center"><iframe width="560" height="315" src="https://www.youtube.com/embed/7WN3N1hnCbI?modestbranding=1" title="YouTube video player" frameborder="0" allow="accelerometer; clipboard-write; encrypted-media; gyroscope; picture-in-picture" allowfullscreen></iframe></p>

projects/rocprofiler-systems/docs/what-is-rocprof-sys.rst

@@ -0,0 +1,33 @@
+.. meta::
+   :description: ROCm Systems Profiler introduction, explanation, and reference
+   :keywords: rocprof-sys, rocprofiler-systems, Omnitrace, ROCm, profiler, explanation, introduction, what is, tracking, visualization, tool, Instinct, accelerator, AMD
+
+******************************
+What is ROCm Systems Profiler?
+******************************
+
+ROCm Systems Profiler is designed for the high-level profiling and comprehensive tracing
+of applications running on the CPU or the CPU and GPU. It supports dynamic binary
+instrumentation, call-stack sampling, and various other features for determining
+which function and line number are currently executing.
+
+A visualization of the comprehensive ROCm Systems Profiler results can be observed in any modern
+web browser. Upload the Perfetto (``.proto``) output files produced by ROCm Systems Profiler at
+`ui.perfetto.dev <https://ui.perfetto.dev/>`_ to see the details.
+
+.. important::
+   If you are using a version of ROCm prior to ROCm 6.3.1 and are experiencing problems viewing your
+   trace in the latest version of [Perfetto](http://ui.perfetto.dev), then try using
+   [Perfetto UI v46.0](https://ui.perfetto.dev/v46.0-35b3d9845/#!/).
+
+Aggregated high-level results are available as human-readable text files and
+JSON files for programmatic analysis. The JSON output files are compatible with the
+`hatchet <https://github.com/hatchet/hatchet>`_ Python package. Hatchet converts
+the performance data into pandas data frames and facilitates multi-run comparisons, filtering,
+and visualization in Jupyter notebooks.
+
+To use ROCm Systems Profiler for instrumentation, follow these two configuration steps:
+
+#. Indicate the functions and modules to :doc:`instrument <./how-to/instrumenting-rewriting-binary-application>` in the target binaries, including the executable and any libraries
+#. Specify the :doc:`instrumentation parameters <./how-to/configuring-runtime-options>` to use when the instrumented binaries are launched
+