diff --git a/projects/rocprofiler/.azuredevops/rocm-ci.yml b/projects/rocprofiler/.azuredevops/rocm-ci.yml new file mode 100644 index 0000000000..27fd51a9ab --- /dev/null +++ b/projects/rocprofiler/.azuredevops/rocm-ci.yml @@ -0,0 +1,40 @@ +resources: + repositories: + - repository: pipelines_repo + type: github + endpoint: ROCm + name: ROCm/ROCm + +variables: +- group: common +- template: /.azuredevops/variables-global.yml@pipelines_repo + +trigger: + batch: true + branches: + include: + - amd-master + - amd-staging + paths: + exclude: + - .github + - doc + - LICENSE + - '*.md' + +pr: + autoCancel: true + branches: + include: + - amd-master + - amd-staging + paths: + exclude: + - .github + - doc + - LICENSE + - '*.md' + drafts: false + +jobs: + - template: ${{ variables.CI_COMPONENT_PATH }}/rocprofiler.yml@pipelines_repo diff --git a/projects/rocprofiler/.clang-format b/projects/rocprofiler/.clang-format new file mode 100644 index 0000000000..7f49cb07ec --- /dev/null +++ b/projects/rocprofiler/.clang-format @@ -0,0 +1,60 @@ +--- +Language: Cpp +BasedOnStyle: Google +AccessModifierOffset: -1 +ConstructorInitializerIndentWidth: 4 +AlignEscapedNewlinesLeft: false +AlignTrailingComments: true +AlignConsecutiveAssignments: false +AlignOperands: false +AllowAllParametersOfDeclarationOnNextLine: true +AllowShortBlocksOnASingleLine: false +AllowShortIfStatementsOnASingleLine: true +AllowShortLoopsOnASingleLine: true +AllowShortFunctionsOnASingleLine: All +AlwaysBreakAfterDefinitionReturnType: false +AlwaysBreakTemplateDeclarations: false +AlwaysBreakBeforeMultilineStrings: true +BreakBeforeBinaryOperators: false +BreakBeforeTernaryOperators: true +BreakConstructorInitializersBeforeComma: false +BinPackParameters: true +ColumnLimit: 100 +ConstructorInitializerAllOnOneLineOrOnePerLine: true +ExperimentalAutoDetectBinPacking: false +IndentCaseLabels: true +IndentWrappedFunctionNames: false +IndentFunctionDeclarationAfterType: false +MaxEmptyLinesToKeep: 2 +KeepEmptyLinesAtTheStartOfBlocks: false +NamespaceIndentation: None +ObjCSpaceAfterProperty: false +ObjCSpaceBeforeProtocolList: false +PenaltyBreakBeforeFirstCallParameter: 1 +PenaltyBreakComment: 300 +PenaltyBreakString: 1000 +PenaltyBreakFirstLessLess: 120 +PenaltyExcessCharacter: 1000000 +PenaltyReturnTypeOnItsOwnLine: 200 +DerivePointerAlignment: false +PointerAlignment: Left +SpacesBeforeTrailingComments: 2 +Cpp11BracedListStyle: true +Standard: Auto +IndentWidth: 2 +TabWidth: 8 +UseTab: Never +BreakBeforeBraces: Attach +SpacesInParentheses: false +SpacesInAngles: false +SpaceInEmptyParentheses: false +SpacesInCStyleCastParentheses: false +SpacesInContainerLiterals: true +SpaceBeforeAssignmentOperators: true +ContinuationIndentWidth: 4 +CommentPragmas: '^ IWYU pragma:' +ForEachMacros: [ foreach, Q_FOREACH, BOOST_FOREACH ] +SpaceBeforeParens: ControlStatements +DisableFormat: false +SortIncludes: false +... diff --git a/projects/rocprofiler/.clang-tidy b/projects/rocprofiler/.clang-tidy new file mode 100644 index 0000000000..c2c3ca3aa9 --- /dev/null +++ b/projects/rocprofiler/.clang-tidy @@ -0,0 +1,30 @@ +--- +Checks: "-*,\ +misc-*,\ +-misc-incorrect-roundings,\ +-misc-macro-parentheses,\ +-misc-misplaced-widening-cast,\ +-misc-static-assert,\ +-misc-no-recursion,\ +-misc-non-private-member-variables-in-classes,\ +modernize-*,\ +-modernize-deprecated-headers,\ +-modernize-raw-string-literal,\ +-modernize-return-braced-init-list,\ +-modernize-use-transparent-functors,\ +-modernize-use-trailing-return-type,\ +-modernize-avoid-c-arrays,\ +-modernize-use-auto,\ +-modernize-concat-nested-namespaces,\ +-modernize-use-nodiscard,\ +performance-*,\ +readability-*,\ +-readability-function-size,\ +-readability-identifier-naming,\ +" +CheckOptions: + - key: readability-braces-around-statements.ShortStatementLines + value: '2' + - key: readability-implicit-bool-conversion.AllowPointerConditions + value: '1' +... diff --git a/projects/rocprofiler/.cmake-format.yaml b/projects/rocprofiler/.cmake-format.yaml new file mode 100644 index 0000000000..1d1aba5305 --- /dev/null +++ b/projects/rocprofiler/.cmake-format.yaml @@ -0,0 +1,98 @@ +parse: + additional_commands: + find_package: + flags: + - EXACT + - QUIET + - MODULE + - REQUIRED + - CONFIG + - NO_MODULE + - GLOBAL + - NO_POLICY_SCOPE + - BYPASS_PROVIDER + - NO_DEFAULT_PATH + - NO_PACKAGE_ROOT_PATH + - NO_CMAKE_PATH + - NO_CMAKE_ENVIRONMENT_PATH + - NO_SYSTEM_ENVIRONMENT_PATH + - NO_CMAKE_PACKAGE_REGISTRY + - NO_CMAKE_BUILDS_PATH + - NO_CMAKE_SYSTEM_PATH + - NO_CMAKE_INSTALL_PREFIX + - NO_CMAKE_SYSTEM_PACKAGE_REGISTRY + - CMAKE_FIND_ROOT_PATH_BOTH + - ONLY_CMAKE_FIND_ROOT_PATH + - NO_CMAKE_FIND_ROOT_PATH + kwargs: + COMPONENTS: '*' + OPTIONAL_COMPONENTS: '*' + NAMES: '*' + CONFIGS: '*' + HINTS: '*' + PATHS: '*' + REGISTRY_VIEW: '*' + PATH_SUFFIXES: '*' + override_spec: {} + vartags: [] + proptags: [] +format: + disable: false + line_width: 90 + tab_size: 4 + use_tabchars: false + fractional_tab_policy: use-space + max_subgroups_hwrap: 2 + max_pargs_hwrap: 8 + max_rows_cmdline: 2 + separate_ctrl_name_with_space: false + separate_fn_name_with_space: false + dangle_parens: false + dangle_align: child + min_prefix_chars: 4 + max_prefix_chars: 10 + max_lines_hwrap: 2 + line_ending: unix + command_case: lower + keyword_case: upper + always_wrap: [] + enable_sort: true + autosort: false + require_valid_layout: false + layout_passes: {} +markup: + bullet_char: '*' + enum_char: . + first_comment_is_literal: true + literal_comment_pattern: ^# + fence_pattern: ^\s*([`~]{3}[`~]*)(.*)$ + ruler_pattern: ^\s*[^\w\s]{3}.*[^\w\s]{3}$ + explicit_trailing_pattern: '#<' + hashruler_min_length: 10 + canonicalize_hashrulers: true + enable_markup: true +lint: + disabled_codes: [] + function_pattern: '[0-9a-z_]+' + macro_pattern: '[0-9A-Z_]+' + global_var_pattern: '[A-Z][0-9A-Z_]+' + internal_var_pattern: _[A-Z][0-9A-Z_]+ + local_var_pattern: '[a-z][a-z0-9_]+' + private_var_pattern: _[0-9a-z_]+ + public_var_pattern: '[A-Z][0-9A-Z_]+' + argument_var_pattern: '[a-z][a-z0-9_]+' + keyword_pattern: '[A-Z][0-9A-Z_]+' + max_conditionals_custom_parser: 2 + min_statement_spacing: 1 + max_statement_spacing: 2 + max_returns: 6 + max_branches: 12 + max_arguments: 5 + max_localvars: 15 + max_statements: 50 +encode: + emit_byteorder_mark: false + input_encoding: utf-8 + output_encoding: utf-8 +misc: + per_command: {} diff --git a/projects/rocprofiler/.github/palamida.yml b/projects/rocprofiler/.github/palamida.yml new file mode 100644 index 0000000000..2a54378973 --- /dev/null +++ b/projects/rocprofiler/.github/palamida.yml @@ -0,0 +1,5 @@ +disabled: false +scmId: gh-emu-rocm +branchesToScan: + - amd-staging + - amd-mainline diff --git a/projects/rocprofiler/.github/workflows/kws-caller.yml b/projects/rocprofiler/.github/workflows/kws-caller.yml new file mode 100644 index 0000000000..c0f4f26807 --- /dev/null +++ b/projects/rocprofiler/.github/workflows/kws-caller.yml @@ -0,0 +1,15 @@ +name: Rocm Validation Suite KWS +on: + push: + branches: [amd-staging, amd-mainline] + pull_request: + types: [opened, synchronize, reopened] + workflow_dispatch: +jobs: + kws: + if: ${{ github.event_name == 'pull_request' }} + uses: AMD-ROCm-Internal/rocm_ci_infra/.github/workflows/kws.yml@mainline + secrets: inherit + with: + pr_number: ${{github.event.pull_request.number}} + base_branch: ${{github.base_ref}} diff --git a/projects/rocprofiler/.github/workflows/rocm-ci-caller.yml b/projects/rocprofiler/.github/workflows/rocm-ci-caller.yml new file mode 100644 index 0000000000..6e11b1282d --- /dev/null +++ b/projects/rocprofiler/.github/workflows/rocm-ci-caller.yml @@ -0,0 +1,25 @@ +name: ROCm CI Caller + +on: + pull_request: + branches: [amd-staging, amd-npi, release/rocm-rel-*] + types: [opened, reopened, synchronize] + push: + branches: [amd-mainline] + workflow_dispatch: + issue_comment: + types: [created] + +jobs: + call-workflow: + if: ${{ github.event_name != 'issue_comment' || github.event.comment.body == '!verify' }} + uses: AMD-ROCm-Internal/rocm_ci_infra/.github/workflows/rocm_ci.yml@mainline + secrets: inherit + with: + input_sha: ${{ github.event_name == 'pull_request' && github.event.pull_request.head.sha || github.sha }} + input_pr_num: ${{ github.event_name == 'pull_request' && github.event.pull_request.number || 0 }} + input_pr_url: ${{ github.event_name == 'pull_request' && github.event.pull_request.html_url || '' }} + input_pr_title: ${{ github.event_name == 'pull_request' && github.event.pull_request.title || '' }} + repository_name: ${{ github.repository }} + base_ref: ${{ github.event_name == 'pull_request' && github.base_ref || github.ref }} + trigger_event_type: ${{ github.event_name }} diff --git a/projects/rocprofiler/.github/workflows/sync-mainline.yaml b/projects/rocprofiler/.github/workflows/sync-mainline.yaml new file mode 100644 index 0000000000..ac4c09af99 --- /dev/null +++ b/projects/rocprofiler/.github/workflows/sync-mainline.yaml @@ -0,0 +1,17 @@ +name: Sync amd-mainline to public repository + +on: + push: + branches: [ amd-mainline ] + +jobs: + git-mirror: + runs-on: ubuntu-latest + steps: + - name: git-sync + uses: AMD-ROCm-Internal/rocprofiler-github-actions@git-sync-v3 + with: + source_repo: "https://${{ secrets.TOKEN }}@github.com/AMD-ROCm-Internal/rocprofiler.git" + source_branch: "amd-mainline" + destination_repo: "https://${{ secrets.EXT_TOKEN }}@github.com/ROCm/rocprofiler.git" + destination_branch: "amd-mainline" diff --git a/projects/rocprofiler/.github/workflows/sync-staging.yaml b/projects/rocprofiler/.github/workflows/sync-staging.yaml new file mode 100644 index 0000000000..c7e82e12ff --- /dev/null +++ b/projects/rocprofiler/.github/workflows/sync-staging.yaml @@ -0,0 +1,17 @@ +name: Sync amd-staging to public repository + +on: + push: + branches: [ amd-staging ] + +jobs: + git-mirror: + runs-on: ubuntu-latest + steps: + - name: git-sync + uses: AMD-ROCm-Internal/rocprofiler-github-actions@git-sync-v3 + with: + source_repo: "https://${{ secrets.TOKEN }}@github.com/AMD-ROCm-Internal/rocprofiler.git" + source_branch: "amd-staging" + destination_repo: "https://${{ secrets.EXT_TOKEN }}@github.com/ROCm/rocprofiler.git" + destination_branch: "amd-staging" diff --git a/projects/rocprofiler/.gitignore b/projects/rocprofiler/.gitignore new file mode 100644 index 0000000000..4e36f8d28b --- /dev/null +++ b/projects/rocprofiler/.gitignore @@ -0,0 +1,6 @@ +build +compile_commands.json +.cache +.DS_Store +plugin/att/__pycache__ +.github diff --git a/projects/rocprofiler/.gitmodules b/projects/rocprofiler/.gitmodules new file mode 100644 index 0000000000..2328b58479 --- /dev/null +++ b/projects/rocprofiler/.gitmodules @@ -0,0 +1,6 @@ + + +[submodule "perfetto"] + path = plugin/perfetto/perfetto + url = https://github.com/google/perfetto.git + branch = releases/v44.x diff --git a/projects/rocprofiler/.markdown-lint-config.yml b/projects/rocprofiler/.markdown-lint-config.yml new file mode 100644 index 0000000000..3161169cb5 --- /dev/null +++ b/projects/rocprofiler/.markdown-lint-config.yml @@ -0,0 +1,141 @@ +default: false # includes/excludes all rules by default + +# Heading levels should only increment by one level at a time +MD001: true + +# Heading style +MD003: true + +# Unordered list style +MD004: true + +# Inconsistent indentation for list items at the same level +MD005: true + +# Consider starting bulleted lists at the beginning of the line +MD006: true + +# Unordered list indentation +MD007: true + +# Trailing spaces +MD009: true + +# Hard tabs +MD010: true + +# Reversed link syntax +MD011: true + +# Multiple consecutive blank lines +MD012: true + +# Line length +MD013: false + +# Dollar signs used before commands without showing output +MD014: false + +# No space after hash on atx style heading +MD018: true + +# Multiple spaces after hash on atx style heading +MD019: true + +# No space inside hashes on closed atx style heading +MD020: true + +# Multiple spaces inside hashes on closed atx style heading +MD021: true + +# Headings should be surrounded by blank lines +MD022: true + +# Headings must start at the beginning of the line +MD023: true + +# Multiple headings with the same content +MD024: + allow_different_nesting: true + +# Multiple top level headings in the same document +MD025: true + +# Trailing punctuation in heading +MD026: true + +# Multiple spaces after blockquote symbol +MD027: true + +# Blank line inside blockquote +MD028: false + +# Ordered list item prefix +MD029: + style: 'one' + +# Spaces after list markers +MD030: true + +# Fenced code blocks should be surrounded by blank lines +MD031: true + +# Lists should be surrounded by blank lines +MD032: true + +# Inline HTML +MD033: true + +# Bare URL used +MD034: true + +# Horizontal rule style +MD035: + style: '***' + +# Emphasis used instead of a heading +MD036: true + +# Spaces inside emphasis markers +MD037: true + +# Spaces inside code span elements +MD038: true + +# Spaces inside link text +MD039: true + +# Fenced code blocks should have a language specified +MD040: true + +# First line in file should be a top level heading +MD041: true + +# No empty links +MD042: true + +# Required heading structure +MD043: false + +# Proper names should have the correct capitalization +MD044: false + +# Images should have alternate text (alt text) +MD045: false + +# Code block style +MD046: + style: 'fenced' + +# Files should end with a single newline character +MD047: true + +# Code fence style +MD048: + style: 'backtick' + +# Custom rules: +CHANGELOG-RULE-001: true +CHANGELOG-RULE-002: true +CHANGELOG-RULE-003: true +CHANGELOG-RULE-004: true \ No newline at end of file diff --git a/projects/rocprofiler/CHANGELOG.md b/projects/rocprofiler/CHANGELOG.md new file mode 100644 index 0000000000..6e418eaee9 --- /dev/null +++ b/projects/rocprofiler/CHANGELOG.md @@ -0,0 +1,395 @@ +# Changelog for ROCprofiler + +Full documentation for ROCprofiler is available at +[docs.amd.com](https://docs.amd.com/bundle/ROCm-Profiling-Tools-User-Guide-v5.3) + +As of ROCm 5.5, the ROCm Profiler will not use terminologies like `rocmtools` or +`rocsight` to describe `rocrofiler` as was done in ROCm 5.4. To identify the +separation of the two versions of `rocprofiler`, the terms `rocprofilerV1` and +`rocprofilerV2` will be used. The `rocprofilerV2` API is currently considered a +beta release and subject to changes in future releases. + +## ROCprofiler for rocm 5.4.4 + +In ROCm 5.4 the naming of the ROCm Profiler related files is: + + | ROCm 5.4 | rocprofilerv1 | rocmtools | + |-----------------|-------------------------------------|---------------------------------| + | **Tool script** | `bin/rocprof` | `bin/rocsight` | + | **API include** | `include/rocprofiler/rocprofiler.h` | `include/rocmtools/rocmtools.h` | + | **API library** | `lib/librocprofiler64.so.1` | `lib/librocmtools.so.1` | + +The ROCm Profiler Tool that uses `rocprofilerV1` can be invoked using the +following command: + +```sh +rocprof … +``` + +To write a custom tool based on the `rocprofilerV1` API do the following: + +```C +main.c: +#include // Use the rocprofilerV1 API +int main() { + // Use the rocprofilerV1 API + return 0; +} +``` + +This can be built in the following manner: + +```sh +gcc main.c -I/opt/rocm-5.4.4/include -L/opt/rocm-5.4.4/lib -lrocprofiler64 +``` + +The resulting `a.out` will depend on +`/opt/rocm-5.4.4/lib/librocprofiler64.so.1`. + +The ROCm Profiler that uses `rocprofilerV2` API can be invoked using the +following command: + +```sh +rocsight … +``` + +To write a custom tool based on the `rocmtools` API do the following: + +```C +main.c: +#include // Use the rocmtools API +int main() { + // Use the rocmtools API + return 0; +} +``` + +This can be built in the following manner: + +```sh +gcc main.c -I/opt/rocm-5.4.4/include -L/opt/rocm-5.4.4/lib -lrocmtools +``` + +The resulting `a.out` will depend on `/opt/rocm-5.4.4/lib/librocmtools.so.1`. + +## ROCprofiler for rocm 5.5.0 + +In ROCm 5.5 the `rocprofilerv1` and `rocprofilerv2` include and library files +are merged into single files. The `rocmtools` available in ROCm 5.4 is also +available in ROCm 5.5 but is deprecated and will be removed in a future release. + + | ROCm 5.5 | rocprofilerv1 | rocprofilerv2 | rocmtools *(deprecated)* | + |-----------------|-------------------------------------|-------------------------------------|---------------------------------| + | **Tool script** | `bin/rocprof` | `bin/rocprofv2` | `bin/rocsight` | + | **API include** | `include/rocprofiler/rocprofiler.h` | `include/rocprofiler/rocprofiler.h` | `include/rocmtools/rocmtools.h` | + | **API library** | `lib/librocprofiler64.so.1` | `lib/librocprofiler64.so.1` | `lib/librocmtools.so.1` | + +The ROCm Profiler Tool that uses `rocprofilerV1` can be invoked using the +following command: + +```sh +rocprof … +``` + +To write a custom tool based on the `rocprofilerV1` API it is necessary to +define the macro `ROCPROFILER_V1`: + +```C +main.c: +#define ROCPROFILER_V1 +#include +int main() { + // Use the rocprofilerV1 API + return 0; +} +``` + +This can be built in the following manner: + +```sh +gcc main.c -I/opt/rocm-5.5.0/include -L/opt/rocm-5.5.0/lib -lrocprofiler64 +``` + +The resulting `a.out` will depend on +`/opt/rocm-5.5.0/lib/librocprofiler64.so.1`. + +The ROCm Profiler that uses `rocprofilerV2` API can be invoked using the +following command: + +```sh +rocprofv2 … +``` + +To write a custom tool based on the `rocprofilerV2` API do the following: + +```C +main.c: +#include +int main() { + // Use the rocprofilerV2 API + return 0; +} +``` + +This can be built in the following manner: + +```sh +gcc main.c -I/opt/rocm-5.5.0/include -L/opt/rocm-5.5.0/lib -lrocprofiler64 +``` + +The resulting `a.out` will depend on +`/opt/rocm-5.5.0/lib/librocprofiler64.so.1`. + +## ROCprofiler for rocm 5.6.0 + +In ROCm 5.6 the `rocprofilerv1` and `rocprofilerv2` include and library files of +ROCm 5.5 are split into separate files. The `rocmtools` files that were +deprecated in ROCm 5.5 have been removed. + + | ROCm 5.6 | rocprofilerv1 | rocprofilerv2 | + |-----------------|-------------------------------------|----------------------------------------| + | **Tool script** | `bin/rocprof` | `bin/rocprofv2` | + | **API include** | `include/rocprofiler/rocprofiler.h` | `include/rocprofiler/v2/rocprofiler.h` | + | **API library** | `lib/librocprofiler.so.1` | `lib/librocprofiler.so.2` | + +The ROCm Profiler Tool that uses `rocprofilerV1` can be invoked using the +following command: + +```sh +rocprof … +``` + +To write a custom tool based on the `rocprofilerV1` API do the following: + +```C +main.c: +#include // Use the rocprofilerV1 API +int main() { + // Use the rocprofilerV1 API + return 0; +} +``` + +This can be built in the following manner: + +```sh +gcc main.c -I/opt/rocm-5.6.0/include -L/opt/rocm-5.6.0/lib -lrocprofiler64 +``` + +The resulting `a.out` will depend on +`/opt/rocm-5.6.0/lib/librocprofiler64.so.1`. + +The ROCm Profiler that uses `rocprofilerV2` API can be invoked using the +following command: + +```sh +rocprofv2 … +``` + +To write a custom tool based on the `rocprofilerV2` API do the following: + +```C +main.c: +#include // Use the rocprofilerV2 API +int main() { + // Use the rocprofilerV2 API + return 0; +} +``` + +This can be built in the following manner: + +```sh +gcc main.c -I/opt/rocm-5.6.0/include -L/opt/rocm-5.6.0/lib -lrocprofiler64v2 +``` + +The resulting `a.out` will depend on +`/opt/rocm-5.6.0/lib/librocprofiler64.so.2`. + +### Optimized + +- Improved Test Suite + +### Added + +- 'end_time' need to be disabled in roctx_trace.txt +- support for hsa_amd_memory_async_copy_on_engine API function trace + +### Fixed + +- rocprof in ROcm/5.4.0 gpu selector broken. +- rocprof in ROCm/5.4.1 fails to generate kernel info. +- rocprof clobbers LD_PRELOAD. + +## ROCprofiler for rocm 5.7.0 + +### Navi support + +Rocprofiler for ROCm 5.7 added support for counter collection (PMC) and advanced thread tracing (ATT) for Navi21 and Navi31 GPUs. + +- On Navi3x, counter collection requires the GPU to be in a stable power state. See README.md for instructions. HIP RT in ATT not yet supported. + +### Changed + +- ATT analysis will not run by default. For ATT to have the same behavior as 5.5, use --plugin att --mode network +- Kernel Names are now removed from HIP API records, users of the API can get the kernel names from the corresponding HIP Dispatch OPS using the correlation ID, this change was done to optimize and to manage the data copied. +- Removing Replay modes as we discovered that some of them will corrupt the applications' behavior, we will re-add them once we implement the fix for them. + +### Optimized + +- Improved ATT parser performance and file sizes. +- Now profiler autocorrect user input errors for pmc and throws exception for wrong input with this message:"Bad input metric. usage --> pmc: [counter1] [counter2]" + +### Added + +- Every API trace in V2 reported synchronously will have two records, one for Enter phase and for Exit phase +- File Plugin now reports the HSA OPS operation kind as part of the output text +- MI300 counters support for rocprof v1 and v2. +- Support for MI300 XCC modes for rocprof v2. +- MI300 individual XCC counters dumped per-xcc as separate records but with same record-id and kernel dispatch info +- Naming for MPI ranks. Filenames containing "%rank" are replaced by variables "MPI_RANK", "OMPI_COMM_WORLD_RANK" or "MV2_COMM_WORLD_RANK". +- MPI Rank will appear in perfetto track names. +- File plugin has been split to File & CLI plugins, CLI plugin is responsible for showing results on the terminal screen and will be automatically the choice if no -d option given in rocprof, File plugin on the other hand is responsible for writing the output results in files if -d option is given. +- Structure of the results is different for both CLI & File plugin; File plugin will make sure every type of result is in a separate file, starting by specifying the header; CLI plugin will have the records in the old way. +Example for file plugin output: + + ```string + Dispatch_ID,GPU_ID,Queue_ID,Queue_Index,PID,TID,GRD,WGR,LDS,SCR,Arch_VGPR,ACCUM_VGPR,SGPR,Wave_Size,SIG,OBJ,Kernel_Name,Start_Timestamp,End_Timestamp,Correlation_ID,GRBM_COUNT + + 1,4,1,1,1584730,1584730,10,10,0,0,8,0,16,64,140464978048000,1,"helloworld(char*, char*) (.kd)",0,140469300947216,33,12637.000000 + ``` + + ```string + Domain,Function,Kernel_Name,Start_Timestamp,End_Timestamp,Correlation_ID + + HIP_API_DOMAIN,hipGetDeviceProperties,,316678074094190,316678074098929,1 + HIP_API_DOMAIN,hipMalloc,,316678074105702,316678074130851,2 + HIP_API_DOMAIN,hipMalloc,,316678074131382,316678074136111,3 + ``` + +- Removing Record IDs from tracer records in CLI plugin. +- Added Flush Interval and Trace Period functionality, where --flush-interval [time_in_ms], for flushing the buffers every given interval by the user, and --trace-period [delay]:[trace_time]:[interval], where delay is the time to wait before starting session, trace_time is the time between every start and stop session and interval the time between two consecutive sessions (omitting interval = infinite). For more details please refer to the ROCProfV2 tool usage document. +- Added requirements.txt to be used to install all the necessary python3 packages. +- ATT plugin: + - Added --mode, --mpi and --depth parameters. + - Limiting file name sizes for large kernels. + - SE_MASK parameter for input.txt, a binary mask specifying for which shader engines to collect from. + On GFX9, SEs are masked out completely. On Navi only part of the data is masked. + The use of SE_MASK=0x1 is heavily encouraged to avoid packet lost events. + - "ROCPROFILER_MAX_ATT_PROFILES" environment variable can be set. Previously fixed at 16, now the default is 1. + - Increased default ATT buffer size per collection to 1GB. Added "BUFFER_SIZE=size" (in MB) parameter to set buffer size. + - Added "DISPATCH=id" or "DISPATCH=id,rank" to set which dispatch ids to profile for which MPI rank. + +### Fixed + +- Samples are fixed to show the new usage of phases. +- Plugin option validates the plugin names. +- Fixing rocsys, for rocsys options, rocsys -h can be called +- "--output-file" option ignored when no output folder was specified. +- Perfetto crash when using ROCTX and/or no output file specified. +- Parsing of the getpc, setpc and swappc instructions with registers loaded from scratch space. +- Some browsers caching ATT data from older kernels. +- Navi2x GPUs required the first counter to be GRBM. This is fixed in 5.7. +- If ROCPROFILER_METRICS_PATH environment variable is not set, the counters xml path will be taken from the following path (../libexec/rocprofiler/counters/derived_counters.xml) which is relative to librocprofiler64.so.2.0.0 +- Repeated base metrics were not being properly reused by derived counters. +- Fixed wrong dispatch ID on kernel.txt + +## ROCprofiler for rocm 6.0 + +### Added + +- Updated supported GPU architectures in README with profiler versions +- Automatic ISA dumping for ATT. See README. +- CSV mode for ATT. See README. +- Added option to control kernel name truncation. +- Limit rocprof(v1) srcipt usage to only supported architectures. +- Added Tool versioning to be able to run rocprofv2 using rocprof. See README for more information. +- Added Plugin Versioning way in rocprofv2. See README for more details. +- Added --version in the rocprof and rocprofv2 to be able to see the current rocprof/v2 version along with ROCm version information. +- Extended rocprof(v1) support for MI300. + +## ROCprofiler for rocm 6.1 + +### Added + +- ATT: Continuous mode. The input file now has a DISPATCH_RANGE=begin,end that specifies continuous ATT collection during the application run. +- ATT: Ability to dump raw codeobjects as .out files with ISA_CAPTURE_MODE=[0,1,2] +- Updated README and CMakeLists.txt for rocm-llvm-dev dependency. + +### Removed + +- ATT: Network mode, since it was redundant with file mode. +- ATT: Generated isa.s file, since the functionality was moved from collection to parsing time. + +### Fixed + +- Multiple fixes and optimizations for ATT (Advanced Thread Tracer). +- Fixed format of presenting MI300 counters for individual XCCs. +- Fixed ROCprofiler to match versioning changes in HIP Runtime. +- Fixed plugins race condition. +- Updated metrics to MI300. + +## ROCprofiler for rocm 6.2 + +### Removed + +- pcsampler sample code has been removed due to deprecation from v2. + +## ROCProfiler for ROCm 6.3 + +### Added + +- JSON output plugin for `rocprofv2`. The JSON file matches Google Trace Format making it easy to load on Perfetto, Chrome tracing, or Speedscope. For Speedscope, use `--disable-json-data-flows` option as speedscope doesn't work with data flows. +- `--no-serialization` flag to disable kernel serialization when `rocprofv2` is in counter collection mode. This allows `rocprofv2` to avoid deadlock when profiling certain programs in counter collection mode. +- `FP64_ACTIVE` and `ENGINE_ACTIVE` metrics to AMD Instinct MI300 accelerator +- New HIP APIs with struct defined inside union. +- Early checks to confirm the eligibility of ELF file in ATT plugin +- Support for kernel name filtering in `rocprofv2` +- Barrier bit to read and stop packets +- ROCProfiler support for gfx1150 and gfx1151 +- ATT support for gfx12 +- gfx12 support + +### Changed + +- Extended lifetime for proxy queues +- Setting the `trace-start` option for `rocprof` to `off` now disables kernel tracing +- `libpciaccess-dev` functions now load with `dlopen` +- `PcieAccessApi*` api and `void* libpciaccess_handle` are now initialized to `nullptr` + +### Removed + +- Obsolete BSD and GPL licenses +- `libsystemd-dev` from `CMakeLists.txt` + +### Optimized + +- ROCProfiler Performance improved to reduce profiling time for large workloads of counter collection + + + +- Fixed bandwidth measurement in MI300 +- Fixed Perfetto plugin issue of `roctx` trace not getting displayed +- Fixed `--help` for counter collection +- Fixed signal management issues in `queue.cpp` +- Fixed Perfetto tracks for multi-GPU +- Fixed Perfetto plugin usage with `rocsys` +- Fixed incorrect number of columns in the output CSV files for counter collection and kernel tracing +- Fixed the ROCProfiler hang issue when running kernel trace, thread trace, or counter collection on Iree benchmark for AMD Instinct MI300 accelerator +- Fixed build errors thrown during parsing of unions +- Fixed the system hang caused while running `--kernel-trace` with Perfetto for certain applications +- Fixed missing profiler records issue caused while running `--trace-period` +- Fixed the hang issue of `ProfilerAPITest` of `runFeatureTests` on AMD Instinct MI300 accelerator +- Fixed segment fault on Navi32 + +## ROCProfiler for ROCm 6.4 + +### Added + +- README updated for kernel filtration +- Tool deprecation message for rocprofiler v1,v2 +- Adding ops 16,32,64 metrics for rdc + +### Resolved issues + +- Fixed issue where invalid UTF-8 characters in a trace log would crash the program diff --git a/projects/rocprofiler/CMakeLists.txt b/projects/rocprofiler/CMakeLists.txt new file mode 100644 index 0000000000..f9f4f15372 --- /dev/null +++ b/projects/rocprofiler/CMakeLists.txt @@ -0,0 +1,598 @@ +# ############################################################################## +# Copyright (c) 2018 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. +# ############################################################################## + +cmake_minimum_required(VERSION 3.18.0 FATAL_ERROR) + +set(CMAKE_C_FLAGS_COVERAGE_INIT + "-Og -g3 -fno-omit-frame-pointer -fprofile-abs-path -fprofile-arcs -ftest-coverage --coverage" + CACHE STRING "C flags for code coverage builds") +set(CMAKE_CXX_FLAGS_COVERAGE_INIT + "-Og -g3 -fno-omit-frame-pointer -fprofile-abs-path -fprofile-arcs -ftest-coverage --coverage" + CACHE STRING "C++ flags for code coverage builds") + +set(CMAKE_C_FLAGS_COVERAGE + "${CMAKE_C_FLAGS_COVERAGE_INIT}" + CACHE STRING "C flags for code coverage builds") +set(CMAKE_CXX_FLAGS_COVERAGE + "${CMAKE_CXX_FLAGS_COVERAGE_INIT}" + CACHE STRING "C++ flags for code coverage builds") + +# Build is not supported on Windows plaform +if(WIN32) + message(FATAL_ERROR "Windows build is not supported.") +endif() + +# Set module name and project name. +set(ROCPROFILER_NAME "rocprofiler") +set(ROCPROFILER_TARGET "${ROCPROFILER_NAME}64") +set(ROCPROFILER_LIBRARY "lib${ROCPROFILER_TARGET}") +project(rocprofiler VERSION 2.0.0) + +if(${ROCM_PATCH_VERSION}) + set(PROJECT_VERSION_PATCH ${ROCM_PATCH_VERSION}) + set(PROJECT_VERSION "${PROJECT_VERSION_MAJOR}.${PROJECT_VERSION_MINOR}.${PROJECT_VERSION_PATCH}") +endif() + +include(GNUInstallDirs) + +# set default ROCM_PATH +if(NOT DEFINED ROCM_PATH) + set(ROCM_PATH + "/opt/rocm" + CACHE STRING "Default ROCM installation directory") +endif() +set(CMAKE_INSTALL_LIBDIR "lib" CACHE STRING "Library install directory") +set(CMAKE_CXX_STANDARD 17) +set(CMAKE_CXX_STANDARD_REQUIRED ON) +set(THREADS_PREFER_PTHREAD_FLAG ON) +set(CMAKE_BUILD_RPATH + "${PROJECT_BINARY_DIR}:${PROJECT_BINARY_DIR}/${CMAKE_INSTALL_LIBDIR}") +set(CMAKE_BUILD_RPATH_USE_ORIGIN ON) +set(CMAKE_SKIP_BUILD_RPATH OFF) + +# Adding default path cmake modules +list(INSERT CMAKE_MODULE_PATH 0 "${CMAKE_CURRENT_SOURCE_DIR}/cmake_modules") + +# Set build environment +include(rocprofiler_options) +include(rocprofiler_utils) +include(rocprofiler_env) +include(rocprofiler_formatting) +include(rocprofiler_linting) + +# Setup the package version. +rocprofiler_get_version("1.0.0") +message("-- LIB-VERSION: ${VERSION_MAJOR}.${VERSION_MINOR}.${VERSION_PATCH}") + +set(BUILD_VERSION_MAJOR ${VERSION_MAJOR}) +set(BUILD_VERSION_MINOR ${VERSION_MINOR}) +set(BUILD_VERSION_PATCH ${VERSION_PATCH}) +if(DEFINED VERSION_BUILD AND NOT ${VERSION_BUILD} STREQUAL "") + message("VERSION BUILD DEFINED ${VERSION_BUILD}") + set(BUILD_VERSION_PATCH "${BUILD_VERSION_PATCH}-${VERSION_BUILD}") +endif() +set(BUILD_VERSION_STRING + "${BUILD_VERSION_MAJOR}.${BUILD_VERSION_MINOR}.${BUILD_VERSION_PATCH}") + +set(LIB_VERSION_MAJOR ${VERSION_MAJOR}) +set(LIB_VERSION_MINOR ${VERSION_MINOR}) +if(${ROCM_PATCH_VERSION}) + set(LIB_VERSION_PATCH ${ROCM_PATCH_VERSION}) +else() + set(LIB_VERSION_PATCH ${VERSION_PATCH}) +endif() +set(LIB_VERSION_STRING "${LIB_VERSION_MAJOR}.${LIB_VERSION_MINOR}.${LIB_VERSION_PATCH}") +message("-- LIB-VERSION STRING: ${LIB_VERSION_STRING}") + +# Set target and root/lib/test directory +set(TARGET_NAME "${ROCPROFILER_TARGET}") +set(ROOT_DIR "${CMAKE_CURRENT_SOURCE_DIR}") +set(LIB_DIR "${ROOT_DIR}/src") +set(TEST_DIR "${ROOT_DIR}/test") + +# To set additional RUNPATH in libraries installed in /opt/rocm-ver/lib/roctracer +set(ROCM_APPEND_PRIVLIB_RPATH "\$ORIGIN/..") + +find_package( + amd_comgr REQUIRED CONFIG + HINTS ${CMAKE_INSTALL_PREFIX} + PATHS ${ROCM_PATH} + PATH_SUFFIXES lib/cmake/amd_comgr) +message(STATUS "Code Object Manager found at ${amd_comgr_DIR}.") +link_libraries(amd_comgr) + +find_package(Threads REQUIRED) +find_package( + hsa-runtime64 REQUIRED CONFIG + HINTS ${CMAKE_INSTALL_PREFIX} + PATHS ${ROCM_PATH}) +find_package( + HIP REQUIRED CONFIG + HINTS ${CMAKE_INSTALL_PREFIX} + PATHS ${ROCM_PATH}) + +find_package( + Clang CONFIG + PATHS "${ROCM_PATH}" + PATH_SUFFIXES "llvm/lib/cmake/clang") + +if(NOT Clang_FOUND) + message( + FATAL_ERROR + "\ +The \"rocm-llvm-dev\" package is not installed. \ +Please install it for your OS distribution. e.g: for Ubuntu: \"sudo apt-get install rocm-llvm-dev\".\ +") +endif() + +find_library(NUMA NAME numa REQUIRED) +link_libraries(${NUMA}) + +get_property( + HSA_RUNTIME_INCLUDE_DIRECTORIES + TARGET hsa-runtime64::hsa-runtime64 + PROPERTY INTERFACE_INCLUDE_DIRECTORIES) +find_file( + HSA_H hsa.h + PATHS ${HSA_RUNTIME_INCLUDE_DIRECTORIES} + PATH_SUFFIXES hsa + NO_DEFAULT_PATH REQUIRED) +get_filename_component(HSA_RUNTIME_INC_PATH ${HSA_H} DIRECTORY) +include_directories(${HSA_RUNTIME_INC_PATH}) + +if(NOT DEFINED LIBRARY_TYPE) + set(LIBRARY_TYPE SHARED) +endif() + +# Enable tracing API +if(NOT USE_PROF_API) + set(USE_PROF_API 1) +endif() + +# Enable use of getROCmInstallPath() API +if(USE_GET_ROCM_PATH_API) + set(USE_GET_ROCM_PATH_API 1) + message(STATUS "Enabled Use of getROCmInstallPath() API" ) +endif() + +# Protocol header lookup +set(PROF_API_HEADER_NAME prof_protocol.h) +if(USE_PROF_API EQUAL 1) + find_path( + PROF_API_HEADER_DIR ${PROF_API_HEADER_NAME} + HINTS ${PROF_API_HEADER_PATH} + PATHS /opt/rocm/include + PATH_SUFFIXES roctracer/ext) + if(NOT PROF_API_HEADER_DIR) + message( + FATAL_ERROR + "Profiling API header not found. Tracer integration disabled. Use -DPROF_API_HEADER_PATH=" + ) + else() + include_directories(${PROF_API_HEADER_DIR}) + message(STATUS "Profiling API: ${PROF_API_HEADER_DIR}/${PROF_API_HEADER_NAME}") + endif() +endif() + +enable_testing() + +# Temporarily for CI to work +set(ROCPROFILER_BUILD_TESTS ON) +set(ROCPROFILER_BUILD_CI ON) + +if(ROCPROFILER_BUILD_CI) + include(CTest) +endif() + +# Build libraries +add_subdirectory(src) + +# Build Plugins +add_subdirectory(plugin) + +if(ROCPROFILER_BUILD_SAMPLES) + # Build samples + add_subdirectory(samples) +endif() + +if(ROCPROFILER_BUILD_TESTS) + # Build tests + add_subdirectory(test) + add_subdirectory(tests-v2) +endif() + +# Installation and packaging +set(DEST_NAME ${ROCPROFILER_NAME}) +if(DEFINED CMAKE_INSTALL_PREFIX) + get_filename_component(prefix_name ${CMAKE_INSTALL_PREFIX} NAME) + get_filename_component(prefix_dir ${CMAKE_INSTALL_PREFIX} DIRECTORY) + if(prefix_name STREQUAL ${DEST_NAME}) + set(CMAKE_INSTALL_PREFIX ${prefix_dir}) + endif() +endif() +if(DEFINED CPACK_PACKAGING_INSTALL_PREFIX) + get_filename_component(prefix_name ${CPACK_PACKAGING_INSTALL_PREFIX} NAME) + get_filename_component(prefix_dir ${CPACK_PACKAGING_INSTALL_PREFIX} DIRECTORY) + if(prefix_name STREQUAL ${DEST_NAME}) + set(CPACK_PACKAGING_INSTALL_PREFIX ${prefix_dir}) + endif() +else() + set(CPACK_PACKAGING_INSTALL_PREFIX ${CMAKE_INSTALL_PREFIX}) +endif() + +set(ORIGINAL_SCRIPT_PATH ${CMAKE_CURRENT_SOURCE_DIR}/bin/tblextr.py) +set(OUTPUT_SCRIPT_PATH ${CMAKE_CURRENT_BINARY_DIR}/bin/tblextr.py) +configure_file(${ORIGINAL_SCRIPT_PATH} ${OUTPUT_SCRIPT_PATH} @ONLY) + +message("CMake-install-prefix: ${CMAKE_INSTALL_PREFIX}") +message("CPack-install-prefix: ${CPACK_PACKAGING_INSTALL_PREFIX}") +message("-----------Dest-name: ${DEST_NAME}") + +# Install headers +install( + FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/core/activity.h + DESTINATION ${CMAKE_INSTALL_INCLUDEDIR}/${ROCPROFILER_NAME} + COMPONENT dev) + +# rpl_run.sh +install( + PROGRAMS ${CMAKE_CURRENT_SOURCE_DIR}/bin/rpl_run.sh + DESTINATION ${CMAKE_INSTALL_BINDIR} + RENAME rocprof + COMPONENT runtime) + +configure_file(bin/rocprofv2 ${PROJECT_BINARY_DIR} COPYONLY) +install( + PROGRAMS ${PROJECT_SOURCE_DIR}/bin/rocprofv2 + DESTINATION ${CMAKE_INSTALL_BINDIR} + COMPONENT runtime) + +install( + PROGRAMS ${CMAKE_CURRENT_SOURCE_DIR}/bin/txt2xml.sh + ${CMAKE_CURRENT_SOURCE_DIR}/bin/merge_traces.sh + DESTINATION ${CMAKE_INSTALL_LIBEXECDIR}/${ROCPROFILER_NAME} + COMPONENT runtime) + +install( + FILES ${CMAKE_CURRENT_SOURCE_DIR}/bin/txt2params.py + ${CMAKE_CURRENT_SOURCE_DIR}/bin/txt2params.py + ${CMAKE_CURRENT_BINARY_DIR}/bin/tblextr.py + ${CMAKE_CURRENT_SOURCE_DIR}/bin/dform.py + ${CMAKE_CURRENT_SOURCE_DIR}/bin/mem_manager.py + ${CMAKE_CURRENT_SOURCE_DIR}/bin/sqlitedb.py + DESTINATION ${CMAKE_INSTALL_LIBEXECDIR}/${ROCPROFILER_NAME} + COMPONENT runtime) + +# gfx_metrics.xml metrics.xml +install( + FILES ${CMAKE_CURRENT_SOURCE_DIR}/test/tool/metrics.xml + ${CMAKE_CURRENT_SOURCE_DIR}/test/tool/gfx_metrics.xml + DESTINATION ${CMAKE_INSTALL_LIBDIR}/${ROCPROFILER_NAME} + COMPONENT runtime) + +if(${LIBRARY_TYPE} STREQUAL SHARED) + # Packaging directives + set(CPACK_GENERATOR "DEB" "RPM" "TGZ" CACHE STRING "CPACK GENERATOR DEB;RPM") + set(ENABLE_LDCONFIG + ON + CACHE BOOL "Set library links and caches using ldconfig.") + set(CPACK_PACKAGE_NAME "${PROJECT_NAME}") + set(CPACK_PACKAGE_VENDOR "Advanced Micro Devices, Inc.") + set(CPACK_PACKAGE_VERSION_MAJOR ${PROJECT_VERSION_MAJOR}) + set(CPACK_PACKAGE_VERSION_MINOR ${PROJECT_VERSION_MINOR}) + set(CPACK_PACKAGE_VERSION_PATCH ${PROJECT_VERSION_PATCH}) + set(CPACK_PACKAGE_VERSION + "${CPACK_PACKAGE_VERSION_MAJOR}.${CPACK_PACKAGE_VERSION_MINOR}.${CPACK_PACKAGE_VERSION_PATCH}" + ) + set(CPACK_PACKAGE_CONTACT "ROCm Profiler Support ") + set(CPACK_PACKAGE_DESCRIPTION_SUMMARY + "ROCPROFILER library for AMD HSA runtime API extension support") + set(CPACK_RESOURCE_FILE_LICENSE "${CMAKE_CURRENT_SOURCE_DIR}/LICENSE") + + if(DEFINED ENV{ROCM_LIBPATCH_VERSION}) + set(CPACK_PACKAGE_VERSION "${CPACK_PACKAGE_VERSION}.$ENV{ROCM_LIBPATCH_VERSION}") + message("Using CPACK_PACKAGE_VERSION ${CPACK_PACKAGE_VERSION}") + endif() + + # Debian package specific variable for ASAN + set(CPACK_DEBIAN_ASAN_PACKAGE_NAME "${ROCPROFILER_NAME}-asan") + set(CPACK_DEBIAN_ASAN_PACKAGE_DEPENDS "hsa-rocr-asan, rocm-core-asan") + + # Install license file + install( + FILES ${CPACK_RESOURCE_FILE_LICENSE} + DESTINATION ${CMAKE_INSTALL_DOCDIR} + COMPONENT runtime) + install( + FILES ${CPACK_RESOURCE_FILE_LICENSE} + DESTINATION ${CMAKE_INSTALL_DOCDIR}-asan + COMPONENT asan) + + # Debian package specific variables + if(DEFINED ENV{CPACK_DEBIAN_PACKAGE_RELEASE}) + set(CPACK_DEBIAN_PACKAGE_RELEASE $ENV{CPACK_DEBIAN_PACKAGE_RELEASE}) + else() + set(CPACK_DEBIAN_PACKAGE_RELEASE "local") + endif() + + message("Using CPACK_DEBIAN_PACKAGE_RELEASE ${CPACK_DEBIAN_PACKAGE_RELEASE}") + set(CPACK_DEB_COMPONENT_INSTALL ON) + set(CPACK_DEBIAN_FILE_NAME "DEB-DEFAULT") + set(CPACK_DEBIAN_RUNTIME_PACKAGE_NAME "${PROJECT_NAME}") + set(CPACK_DEBIAN_RUNTIME_PACKAGE_DEPENDS + "rocminfo, hsa-rocr, rocm-core, libelf-dev, libnuma-dev, libxml2-dev" + ) + set(CPACK_DEBIAN_DEV_PACKAGE_NAME "${PROJECT_NAME}-dev") + set(CPACK_DEBIAN_DEV_PACKAGE_DEPENDS "${PROJECT_NAME}, hsa-rocr-dev, rocm-core") + set(CPACK_DEBIAN_TESTS_PACKAGE_NAME "${PROJECT_NAME}-tests") + set(CPACK_DEBIAN_TESTS_PACKAGE_DEPENDS "${PROJECT_NAME}-dev, hsa-rocr-dev, rocm-core") + set(CPACK_DEBIAN_SAMPLES_PACKAGE_NAME "${PROJECT_NAME}-samples") + set(CPACK_DEBIAN_SAMPLES_PACKAGE_DEPENDS + "${PROJECT_NAME}-dev, hsa-rocr-dev, rocm-core") + set(CPACK_DEBIAN_DOCS_PACKAGE_NAME "${PROJECT_NAME}-docs") + set(CPACK_DEBIAN_DOCS_PACKAGE_DEPENDS "${PROJECT_NAME}-dev, hsa-rocr-dev, rocm-core") + set(CPACK_DEBIAN_PLUGINS_PACKAGE_NAME "${PROJECT_NAME}-plugins") + set(CPACK_DEBIAN_PLUGINS_PACKAGE_DEPENDS "${PROJECT_NAME}, hsa-rocr-dev, rocm-core") + + set(CPACK_DEBIAN_CHANGELOG_FILE "${CMAKE_CURRENT_SOURCE_DIR}/CHANGELOG.md") + + # RPM package specific variables + if(DEFINED ENV{CPACK_RPM_PACKAGE_RELEASE}) + set(CPACK_RPM_PACKAGE_RELEASE $ENV{CPACK_RPM_PACKAGE_RELEASE}) + else() + set(CPACK_RPM_PACKAGE_RELEASE "local") + endif() + + message("Using CPACK_RPM_PACKAGE_RELEASE ${CPACK_RPM_PACKAGE_RELEASE}") + + set(CPACK_RPM_PACKAGE_LICENSE "MIT") + + # 'dist' breaks manual builds on debian systems due to empty Provides + execute_process( + COMMAND rpm --eval %{?dist} + RESULT_VARIABLE PROC_RESULT + OUTPUT_VARIABLE EVAL_RESULT + OUTPUT_STRIP_TRAILING_WHITESPACE) + message("RESULT_VARIABLE ${PROC_RESULT} OUTPUT_VARIABLE: ${EVAL_RESULT}") + + if(PROC_RESULT EQUAL "0" AND NOT EVAL_RESULT STREQUAL "") + string(APPEND CPACK_RPM_PACKAGE_RELEASE "%{?dist}") + endif() + + set(CPACK_RPM_COMPONENT_INSTALL ON) + set(CPACK_RPM_FILE_NAME "RPM-DEFAULT") + set(CPACK_RPM_RUNTIME_PACKAGE_NAME "${PROJECT_NAME}") + set(CPACK_RPM_RUNTIME_PACKAGE_REQUIRES + "rocminfo, hsa-rocr, rocm-core, libxml2-devel") + set(CPACK_RPM_DEV_PACKAGE_NAME "${PROJECT_NAME}-devel") + set(CPACK_RPM_DEV_PACKAGE_REQUIRES "${PROJECT_NAME}, hsa-rocr-devel, rocm-core") + set(CPACK_RPM_DEV_PACKAGE_PROVIDES "${PROJECT_NAME}-dev") + set(CPACK_RPM_DEV_PACKAGE_OBSOLETES "${PROJECT_NAME}-dev") + set(CPACK_RPM_TESTS_PACKAGE_NAME "${PROJECT_NAME}-tests") + set(CPACK_RPM_TESTS_PACKAGE_REQUIRES "${PROJECT_NAME}-devel, hsa-rocr-devel, rocm-core") + set(CPACK_RPM_DOCS_PACKAGE_NAME "${PROJECT_NAME}-docs") + set(CPACK_RPM_DOCS_PACKAGE_REQUIRES "${PROJECT_NAME}-devel, hsa-rocr-devel, rocm-core") + set(CPACK_RPM_PLUGINS_PACKAGE_NAME "${PROJECT_NAME}-plugins") + set(CPACK_RPM_PLUGINS_PACKAGE_REQUIRES "${PROJECT_NAME}, hsa-rocr-devel, rocm-core") + set(CPACK_RPM_PACKAGE_AUTOREQ 0) + set(CPACK_RPM_SAMPLES_PACKAGE_NAME "${PROJECT_NAME}-samples") + set(CPACK_RPM_SAMPLES_PACKAGE_REQUIRES + "${PROJECT_NAME}-devel, hsa-rocr-devel, rocm-core, hip-runtime-amd") + message("CPACK_RPM_PACKAGE_RELEASE: ${CPACK_RPM_PACKAGE_RELEASE}") + + # Disable build id for rocprofiler as its creating transaction error + set(CPACK_RPM_SPEC_MORE_DEFINE + "%define _build_id_links none + %global __strip ${CPACK_STRIP_EXECUTABLE} + %global __objdump ${CPACK_OBJDUMP_EXECUTABLE} + %global __objcopy ${CPACK_OBJCOPY_EXECUTABLE} + %global __readelf ${CPACK_READELF_EXECUTABLE}") + + # RPM package specific variable for ASAN + set(CPACK_RPM_ASAN_PACKAGE_NAME "${ROCPROFILER_NAME}-asan") + set(CPACK_RPM_ASAN_PACKAGE_REQUIRES "hsa-rocr-asan, rocm-core-asan") + + # set ( CPACK_RPM_CHANGELOG_FILE "${CMAKE_CURRENT_SOURCE_DIR}/CHANGELOG.md" ) + + # Remove dependency on rocm-core if -DROCM_DEP_ROCMCORE=ON not given to cmake + if(NOT ROCM_DEP_ROCMCORE) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_RPM_RUNTIME_PACKAGE_REQUIRES + ${CPACK_RPM_RUNTIME_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_RPM_DEV_PACKAGE_REQUIRES + ${CPACK_RPM_DEV_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core-asan" "" CPACK_RPM_ASAN_PACKAGE_REQUIRES + ${CPACK_RPM_ASAN_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_RPM_TESTS_PACKAGE_REQUIRES + ${CPACK_RPM_TESTS_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_RPM_SAMPLES_PACKAGE_REQUIRES + ${CPACK_RPM_SAMPLES_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_RPM_DOCS_PACKAGE_REQUIRES + ${CPACK_RPM_DOCS_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_RPM_PLUGINS_PACKAGE_REQUIRES + ${CPACK_RPM_PLUGINS_PACKAGE_REQUIRES}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_DEBIAN_RUNTIME_PACKAGE_DEPENDS + ${CPACK_DEBIAN_RUNTIME_PACKAGE_DEPENDS}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_DEBIAN_DEV_PACKAGE_DEPENDS + ${CPACK_DEBIAN_DEV_PACKAGE_DEPENDS}) + string(REGEX REPLACE ",? ?rocm-core-asan" "" CPACK_DEBIAN_ASAN_PACKAGE_DEPENDS + ${CPACK_DEBIAN_ASAN_PACKAGE_DEPENDS}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_DEBIAN_TESTS_PACKAGE_DEPENDS + ${CPACK_DEBIAN_TESTS_PACKAGE_DEPENDS}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_DEBIAN_SAMPLES_PACKAGE_DEPENDS + ${CPACK_DEBIAN_SAMPLES_PACKAGE_DEPENDS}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_DEBIAN_DOCS_PACKAGE_DEPENDS + ${CPACK_DEBIAN_DOCS_PACKAGE_DEPENDS}) + string(REGEX REPLACE ",? ?rocm-core" "" CPACK_DEBIAN_PLUGINS_PACKAGE_DEPENDS + ${CPACK_DEBIAN_PLUGINS_PACKAGE_DEPENDS}) + endif() + + # set components + if(ENABLE_ASAN_PACKAGING) + # ASAN Package requires only asan component with libraries and license file + set(CPACK_COMPONENTS_ALL asan) + else() + set(CPACK_COMPONENTS_ALL runtime dev tests docs plugins samples) + endif() + + include(CPack) + + cpack_add_component( + runtime + DISPLAY_NAME "Runtime" + DESCRIPTION "Dynamic libraries for the ROCProfiler") + + cpack_add_component( + dev + DISPLAY_NAME "Development" + DESCRIPTION "Development needed header files for ROCProfiler" + DEPENDS runtime) + + cpack_add_component( + plugins + DISPLAY_NAME "ROCProfile Plugins" + DESCRIPTION "Plugins for handling ROCProfiler data output" + DEPENDS runtime) + + cpack_add_component( + tests + DISPLAY_NAME "Tests" + DESCRIPTION "Tests for the ROCProfiler" + DEPENDS dev) + + cpack_add_component( + samples + DISPLAY_NAME "Samples" + DESCRIPTION "Samples for the ROCProfiler" + DEPENDS dev) + + cpack_add_component( + docs + DISPLAY_NAME "Documentation" + DESCRIPTION "Documentation for the ROCProfiler API" + DEPENDS dev) + + cpack_add_component( + asan + DISPLAY_NAME "ASAN" + DESCRIPTION "ASAN libraries for the ROCPROFILER" + DEPENDS asan) +endif() + +find_package(Doxygen) + +if(DOXYGEN_FOUND) + # # Set input and output files for API Document + set(DOXYGEN_IN ${CMAKE_CURRENT_SOURCE_DIR}/doc/Doxyfile_API.in) + set(DOXYGEN_OUT ${CMAKE_CURRENT_BINARY_DIR}/Doxyfile_API) + + # # Request to configure the file + configure_file(${DOXYGEN_IN} ${DOXYGEN_OUT} @ONLY) + + add_custom_command( + OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/doc/html/index.html + ${CMAKE_CURRENT_BINARY_DIR}/doc/latex/refman.pdf + COMMAND ${DOXYGEN_EXECUTABLE} ${DOXYGEN_OUT} + COMMAND make -C ${CMAKE_CURRENT_BINARY_DIR}/doc/latex pdf + MAIN_DEPENDENCY ${DOXYGEN_OUT} + ${DOXYGEN_IN} + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/include/rocprofiler/v2/rocprofiler_plugin.h + ${CMAKE_CURRENT_SOURCE_DIR}/include/rocprofiler/v2/rocprofiler.h + COMMENT "Generating API documentation") + + add_custom_target(doc DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/doc/html/index.html + ${CMAKE_CURRENT_BINARY_DIR}/doc/latex/refman.pdf) + + install( + FILES "${CMAKE_CURRENT_BINARY_DIR}/doc/latex/refman.pdf" + DESTINATION ${CMAKE_INSTALL_DOCDIR} + RENAME "${PROJECT_NAME}v2_api_spec.pdf" + OPTIONAL + COMPONENT docs) + + install( + DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/doc/html/" + DESTINATION ${CMAKE_INSTALL_DOCDIR}/html + OPTIONAL + COMPONENT docs) + + # # Set input and output files for Tools Document + set(DOXYGEN_IN ${CMAKE_CURRENT_SOURCE_DIR}/doc/Doxyfile_Tool.in) + set(DOXYGEN_OUT ${CMAKE_CURRENT_BINARY_DIR}/tooldoc/Doxyfile_Tool) + + # # Request to configure the file + configure_file(${DOXYGEN_IN} ${DOXYGEN_OUT} @ONLY) + + add_custom_command( + OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/tooldoc/html/index.html + ${CMAKE_CURRENT_BINARY_DIR}/tooldoc/latex/refman.pdf + COMMAND ${DOXYGEN_EXECUTABLE} ${DOXYGEN_OUT} + COMMAND make -C ${CMAKE_CURRENT_BINARY_DIR}/tooldoc/latex pdf + MAIN_DEPENDENCY ${DOXYGEN_OUT} + ${DOXYGEN_IN} + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/doc/rocprofv2_tool.md + COMMENT "Generating Tools documentation") + + add_custom_target( + doc_tool DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/tooldoc/html/index.html + ${CMAKE_CURRENT_BINARY_DIR}/tooldoc/latex/refman.pdf) + + install( + FILES "${CMAKE_CURRENT_BINARY_DIR}/tooldoc/latex/refman.pdf" + DESTINATION ${CMAKE_INSTALL_DOCDIR} + RENAME "${PROJECT_NAME}v2_tool.pdf" + OPTIONAL + COMPONENT docs) + + install( + DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/tooldoc/html/" + DESTINATION ${CMAKE_INSTALL_DOCDIR}/html + OPTIONAL + COMPONENT docs) + + # # Set input and output files for changelog document + set(DOXYGEN_IN ${CMAKE_CURRENT_SOURCE_DIR}/doc/Doxyfile_ChangeLog.in) + set(DOXYGEN_OUT ${CMAKE_CURRENT_BINARY_DIR}/doc/changelog/Doxyfile_ChangeLog) + + # # Request to configure the file + configure_file(${DOXYGEN_IN} ${DOXYGEN_OUT} @ONLY) + + add_custom_command( + OUTPUT ${CMAKE_CURRENT_BINARY_DIR}/doc/changelog/latex/refman.pdf + COMMAND ${DOXYGEN_EXECUTABLE} ${DOXYGEN_OUT} + COMMAND make -C ${CMAKE_CURRENT_BINARY_DIR}/doc/changelog/latex pdf + MAIN_DEPENDENCY ${DOXYGEN_OUT} + ${DOXYGEN_IN} + DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/CHANGELOG.md + COMMENT "Generating changelog documentation") + + add_custom_target(doc_changelog + DEPENDS ${CMAKE_CURRENT_BINARY_DIR}/doc/changelog/latex/refman.pdf) + + install( + FILES "${CMAKE_CURRENT_BINARY_DIR}/doc/changelog/latex/refman.pdf" + DESTINATION ${CMAKE_INSTALL_DOCDIR} + RENAME "${PROJECT_NAME}_ChangeLog.pdf" + OPTIONAL + COMPONENT docs) + + add_dependencies(doc doc_changelog) +endif() diff --git a/projects/rocprofiler/CODEOWNERS b/projects/rocprofiler/CODEOWNERS new file mode 100644 index 0000000000..7c775e6dbe --- /dev/null +++ b/projects/rocprofiler/CODEOWNERS @@ -0,0 +1 @@ +* @ammarwa @bgopesh diff --git a/projects/rocprofiler/DEBIAN/postinst.in b/projects/rocprofiler/DEBIAN/postinst.in new file mode 100644 index 0000000000..33848c4af1 --- /dev/null +++ b/projects/rocprofiler/DEBIAN/postinst.in @@ -0,0 +1,23 @@ +#!/bin/bash + +set -e + +do_ldconfig() { + # left-hand term originates from ENABLE_LDCONFIG = ON/OFF at package build + if [ "@ENABLE_LDCONFIG@" == "ON" ]; then + echo @CPACK_PACKAGING_INSTALL_PREFIX@/@CMAKE_INSTALL_LIBDIR@ > /@CMAKE_INSTALL_SYSCONFDIR@/ld.so.conf.d/librocprofiler64.conf + ldconfig + fi +} + +case "$1" in + ( configure ) + do_ldconfig + ;; + ( abort-upgrade | abort-remove | abort-deconfigure ) + echo "$1" + ;; + ( * ) + exit 0 + ;; +esac diff --git a/projects/rocprofiler/DEBIAN/prerm.in b/projects/rocprofiler/DEBIAN/prerm.in new file mode 100644 index 0000000000..88dc4aab74 --- /dev/null +++ b/projects/rocprofiler/DEBIAN/prerm.in @@ -0,0 +1,22 @@ +#!/bin/bash + +set -e + +rm_ldconfig() { + # left-hand term originates from ENABLE_LDCONFIG = ON/OFF at package build + if [ "@ENABLE_LDCONFIG@" == "ON" ]; then + rm -f /@CMAKE_INSTALL_SYSCONFDIR@/ld.so.conf.d/librocprofiler64.conf + ldconfig + fi +} + +case "$1" in + ( remove | upgrade ) + rm_ldconfig + ;; + ( purge ) + ;; + ( * ) + exit 0 + ;; +esac diff --git a/projects/rocprofiler/LICENSE b/projects/rocprofiler/LICENSE new file mode 100644 index 0000000000..b43445708d --- /dev/null +++ b/projects/rocprofiler/LICENSE @@ -0,0 +1,20 @@ +Copyright (c) 2018-2025 Advanced Micro Devices, Inc. All rights reserved. +[MITx11 License] + +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. \ No newline at end of file diff --git a/projects/rocprofiler/README.md b/projects/rocprofiler/README.md new file mode 100644 index 0000000000..ba5fe33878 --- /dev/null +++ b/projects/rocprofiler/README.md @@ -0,0 +1,598 @@ +# ROCm Profiling Tools + +> [!IMPORTANT] +We are phasing out development and support for roctracer/rocprofiler/rocprof/rocprofv2 in favor of rocprofiler-sdk/rocprofv3 in upcoming ROCm releases. Going forward, only critical defect fixes will be addressed for older versions of profiling tools and libraries. We encourage all users to upgrade to the latest version, rocprofiler-sdk library and rocprofv3 tool, to ensure continued support and access to new features. + +> [!NOTE] +> The published documentation is available at [ROCProfiler documentation](https://rocm.docs.amd.com/projects/rocprofiler/en/latest/index.html) in an organized, easy-to-read format, with search and a table of contents. + +[![Build Status]( +https://dev.azure.com/ROCm-CI/ROCm-CI/_apis/build/status%2Frocprofiler?repoName=ROCm%2Frocprofiler&branchName=amd-staging)](https://dev.azure.com/ROCm-CI/ROCm-CI/_build/latest?definitionId=143&repoName=ROCm%2Frocprofiler&branchName=amd-staging) + +## Introduction + +ROCProfiler is AMD’s tooling infrastructure that provides a hardware specific low level performance analysis interface for the profiling and the tracing of GPU compute applications. + +## ROCProfiler V1 + +Profiling with metrics and traces based on perfcounters (PMC) and traces (SPM). +Implementation is based on AqlProfile HSA extension. +The last API library version for ROCProfiler v1 is 8.0.0 + +The library source tree: + +- doc - Documentation +- include/rocprofiler/rocprofiler.h - Library public API +- include/rocprofiler/v2/rocprofiler.h - V2 Beta Library public API +- include/rocprofiler/v2/rocprofiler_plugins.h - V2 Beta Tool's Plugins Library public API +- src - Library sources + - core - Library API sources + - util - Library utils sources + - xml - XML parser +- test - Library test suite + - ctrl - Test controll + - util - Test utils + - simple_convolution - Simple convolution test kernel + +## Build environment + +Roctracer & Rocprofiler need to be installed in the same directory. + +```bash +export CMAKE_PREFIX_PATH=: +export CMAKE_BUILD_TYPE= # release by default +export CMAKE_DEBUG_TRACE=1 # 1 to enable debug tracing +``` + +To build with the current installed ROCM: + +```bash +cd .../rocprofiler +./build.sh ## (for clean build use `-cb`) +``` + +To run the test: + +```bash +cd .../rocprofiler/build +export LD_LIBRARY_PATH=.:$LD_LIBRARY_PATH # paths to ROC profiler and oher libraries +export HSA_TOOLS_LIB=librocprofiler64.so.1 # ROC profiler library loaded by HSA runtime +export ROCP_TOOL_LIB=test/librocprof-tool.so # tool library loaded by ROC profiler +export LIBPCI_TOOL_LIB=libpciaccess.so # PCI access library loaded by ROC profiler +export ROCP_METRICS=metrics.xml # ROC profiler metrics config file +export ROCP_INPUT=input.xml # input file for the tool library +export ROCP_OUTPUT_DIR=./ # output directory for the tool library, for metrics results file 'results.txt' and trace files +./ +``` + +Internal 'simple_convolution' test run script: + +```bash +cd .../rocprofiler/build +./run.sh +``` + +- To enabled error messages logging to '/tmp/rocprofiler_log.txt': + +```bash +export ROCPROFILER_LOG=1 +``` + +- To enable verbose tracing: + +```bash +export ROCPROFILER_TRACE=1 +``` + +## Supported AMD GPU Architectures (V1) + + The following AMD GPU architectures are supported with ROCprofiler V1: + +- gfx8 (Fiji/Ellesmere) +- gfx900 (AMD Vega 10) +- gfx906 (AMD Vega 7nm also referred to as AMD Vega 20) +- gfx908 (AMD Instinct™ MI100 accelerator) +- gfx90a (AMD Instinct™ MI200) +- gfx94x (AMD Instinct™ MI300) + +*** +Note: ROCProfiler V1 tool usage documentation is available at [Click Here](doc/rocprof_tool.md) +*** + +## ROCProfiler V2 + +The first API library version for ROCProfiler v2 is 9.0.0 + +*** +Note: ROCProfilerV2 is currently considered a beta version and is subject to change in future releases +*** + +### ROCProfilerV2 Modules + +- Counters +- Hardware +- Generic Buffer +- Session +- Filter +- Tools +- Plugins +- Samples +- Tests + +## Getting started + +### Requirements + +- rocm-llvm-dev +- makecache +- Gtest Development Package (Ubuntu: libgtest-dev) +- libelf-dev, libnuma-dev, libpciaccess-dev on ubuntu or their corresponding packages on any other OS +- Cppheaderparser, websockets, matplotlib, lxml, barectf Python3 Packages +- Python packages can be installed using: + + ```bash + pip3 install -r requirements.txt + ``` + +### Build + +The user has two options for building: + +- Option 1 (It will install in the path saved in ROCM_PATH environment variable or /opt/rocm if ROCM_PATH is empty): + + - Run + + Normal Build + + ```bash + ./build.sh --build OR ./build.sh -b + ``` + + Clean Build + + ```bash + ./build.sh --clean-build OR ./build.sh -cb + ``` + +- Option 2 (Where ROCM_PATH environment need to be set with the current installation directory of rocm), run the following: + + - Creating the build directory + + ```bash + mkdir build && cd build + ``` + + - Configuring the rocprofv2 build + + ```bash + cmake -DCMAKE_PREFIX_PATH=$ROCM_PATH -DCMAKE_MODULE_PATH=$ROCM_PATH/hip/cmake -DROCPROFILER_BUILD_TESTS=1 -DROCPROFILER_BUILD_SAMPLES=1 .. + ``` + + - Building the main runtime of the rocprofv2 project + + ```bash + cmake --build . -- -j + ``` + + - Optionally, for building API documentation + + ```bash + cmake --build . -- -j doc + ``` + + - Optionally, for building packages (DEB, RPM, TGZ) + Note: Requires rpm package on ubuntu + + ```bash + cmake --build . -- -j package + ``` + +### Install + +- Optionally, run the following to install + + ```bash + cd build + cmake --build . -- -j install + ``` + +## Features & Usage + +### rocsys + +A command line utility to control a session (launch/start/stop/exit), with the required application to be traced or profiled in a rocprofv2 context. Usage: + +- Launch the application with the required profiling and tracing options with giving a session identifier to be used later + + ```bash + rocsys --session session_name launch mpiexec -n 2 rocprofv2 -i samples/input.txt Histogram + ``` + +- Start a session with a given identifier created at launch + + ```bash + rocsys --session session_name start + ``` + +- Stop a session with a given identifier created at launch + + ```bash + rocsys –session session_name stop + ``` + +- Exit a session with a given identifier created at launch + + ```bash + rocsys –session session_name exit + ``` + +### ROCProf Versioning Support + +Currently, rocprof can support both versions, rocprof and rocprofv2, that can be done using `--tool-version` + +```bash +rocprof --tool-version +``` + +- `--tool-version 1` means it will just use rocprof V1. +- `--tool-version 2` means it will just use rocprofv2. + +To know what version you are using right now, along with more information about the rocm version, use the following: + +```bash +rocprof --version +``` + +### Counters and Metric Collection + +HW counters and derived metrics can be collected using following option: + +```bash +rocprofv2 -i samples/input.txt +``` + +input.txt content Example (Details of what is needed inside input.txt will be mentioned with every feature): + + `pmc: SQ_WAVES GRBM_COUNT GRBM_GUI_ACTIVE SQ_INSTS_VALU` + +By default, rocprofv2 runs counter-collection mode with kernel serialization enabled. With serialization enabled, there is a possibility of deadlock occurring if a program that is being profiled attempts to launch two kernels that rely on each other to progress (e.g. kernel A waits in a while-loop until kernel B modifies a location in memory, but, due to kernel serialization, kernel B will not be launched until kernel A finishes executing). When encountering this issue, kernel serialization can be disabled with the following option: + +```bash +rocprofv2 -ns -i samples/input.txt +``` + +or + +```bash +rocprofv2 --no-serialization -i samples/input.txt +``` + +Alternatively, an environment variable can be set as follows: + +```bash +export ROCPROFILER_NO_SERIALIZATION = 1 +``` + +### Application Trace Support + +Different trace options are available while profiling an app: + +- HIP API & asynchronous activity tracing + + ```bash + rocprofv2 --hip-api ## For synchronous HIP API Activity tracing + rocprofv2 --hip-activity ## For both Synchronous & ASynchronous HIP API Activity tracing + rocprofv2 --hip-trace ## Same as --hip-activity, added for backward compatibility + ``` + +- HSA API & asynchronous activity tracing + + ```bash + rocprofv2 --hsa-api ## For synchronous HSA API Activity tracing + rocprofv2 --hsa-activity ## For both Synchronous & ASynchronous HSA API Activity tracing + rocprofv2 --hsa-trace ## Same as --hsa-activity, added for backward compatibility + ``` + +- Kernel dispatches tracing + + ```bash + rocprofv2 --kernel-trace ## Kernel Dispatch Tracing + ``` + +- HIP & HSA API and asynchronous activity and kernel dispatches tracing + + ```bash + rocprofv2 --sys-trace ## Same as combining --hip-trace & --hsa-trace & ROCtx trace + ``` + +- For complete usage options, please run rocprofv2 help + + ```bash + rocprofv2 --help + ``` + +### Plugin Support + +We have a template for adding new plugins. New plugins can be written on top of rocprofv2 to support the desired output format using include/rocprofiler/v2/rocprofiler_plugins.h header file. These plugins are modular in nature and can easily be decoupled from the code based on need. Installation files: + +```string +rocprofiler-plugins_2.0.0-local_amd64.deb +rocprofiler-plugins-2.0.0-local.x86_64.rpm +``` + +Plugins may have multiple versions, the user can specify which version of the plugin to use by running the following command: + +```bash +rocprofv2 --plugin --plugin-version +``` + +- File plugin: outputs the data in txt files. File plugin have two versions, by default version 2 is the current default. +Usage: + + ```bash + rocprofv2 --plugin file -i samples/input.txt -d output_dir # -d is optional, but can be used to define the directory output for output results + ``` + + File plugin version 1 output header will be similar to the legacy rocprof v1 output: + + ```text + Index,KernelName,gpu-id,queue-id,queue-index,pid,tid,grd,wgr,lds,scr,arch_vgpr,accum_vgpr,sgpr,wave_size,sig,obj,DispatchNs,BeginNs,EndNs,CompleteNs,Counters + ``` + + File plugin version 2 output header: + + ```text + Dispatch_ID,GPU_ID,Queue_ID,PID,TID,Grid_Size,Workgroup_Size,LDS_Per_Workgroup,Scratch_Per_Workitem,Arch_VGPR,Accum_VGPR,SGPR,Wave_Size,Kernel_Name,Start_Timestamp,End_Timestamp,Correlation_ID,Counters + ``` + +- Perfetto plugin: outputs the data in protobuf format. Protobuf files can be viewed using ui.perfetto.dev or using trace_processor. +Usage: + + ```bash + rocprofv2 --plugin perfetto --hsa-trace -d output_dir # -d is optional, but can be used to define the directory output for output results + ``` + +- CTF plugin: Outputs the data in ctf format(a binary trace format). CTF binary output can be viewed using TraceCompass or babeltrace. +Usage: + + ```bash + rocprofv2 --plugin ctf --hip-trace -d output_dir # -d is optional, but can be used to define the directory output for output results + ``` + +- JSON plugin: Outputs `.json` file, the JSON file matches Google Trace Format, so it should be easily loaded to perfetto, chrome tracing or speedscope. For Speedscope, `--disable-json-data-flows` option will be needed as speedscope doesn't work with data flows. +Usage: + + ```bash + rocprofv2 --plugin json --hip-trace -d output_dir + ``` + +### Flush Interval + +Flush interval can be used to control the interval time in milliseconds between the buffers flush for the tool. However, if the buffers are full the flush will be called on its own. This can be used as in the next example: + +```bash +rocprofv2 --flush-interval +``` + +### Trace Period + +Trace period can be used to control when the profiling or tracing is enabled using two arguments, the first one is the delay time, which is the time spent idle without tracing or profiling. The second argument is the profiling or the tracing time, which is the active time where the profiling and tracing are working, so basically, the session will work in the following timeline: + +```string + => => => +``` + + This feature can be used using the following command: + +```bash +rocprofv2 --trace-period :: +``` + +- delay: Time delay to start profiling (ms). +- active_time: How long to profile for (ms). +- interval: If set, profiling sessions will start (loop) every "interval", and run for "active_time", until the application ends. Must be higher than "active_time". + +### Device Profiling + +A device profiling session allows the user to profile the GPU device for counters irrespective of the running applications on the GPU. This is different from application profiling. device profiling session doesn't care about the host running processes and threads. It directly provides low level profiling information. + +### Session Support + + A session is a unique identifier for a profiling/tracing/pc-sampling task. A ROCProfilerV2 Session has enough information about what needs to be collected or traced and it allows the user to start/stop profiling/tracing whenever required. More details on the API can be found in the API specification documentation that can be installed using rocprofiler-doc package. Samples also can be found for how to use the API in samples directory. + +### Kernel Filtration + +rocprofv2 supports the ``kernel`` attribute in input files for only ``csv`` output format. +To enable kernel filtration with input file, the row beginning with the ``kernel:`` keyword should specify the names of kernels to be profiled.Filtering can also be enabled by the flag ``ROCPROFILER_KERNEL_FILTER``. + +For a sample ``vecCopy`` kernel: + +```CPP +__global__ void vecCopy(double *a, double *b, double *c, int n, int stride) { + // Get our global thread ID + int id = blockIdx.x*blockDim.x+threadIdx.x; + if (id < n) + c[id] = a[id]; +} +``` + +To filter vecCopy kernel: + +```bash +$ cat input.txt +pmc: SQ_WAVES +kernel: vecCopy +``` + +```bash +rocprofv2 -i input.txt -d out +``` + +## Tests + + We make use of the GoogleTest (Gtest) framework to automatically find and add test cases to the CMAKE testing environment. ROCProfilerV2 testing is categorized as following: + +- unittests (Gtest Based) : These includes tests for core classes. Any newly added functionality should have a unit test written to it. + +- featuretests (standalone and Gtest Based): These includes both API tests and tool tests. Tool is tested against different applications to make sure we have right output in every run. + +- memorytests (standalone): This includes running address sanitizer for memory leaks, corruptions. + +installation: +rocprofiler-tests_9.0.0-local_amd64.deb +rocprofiler-tests-9.0.0-local.x86_64.rpm + +### List and Run tests + +#### Run unit tests on the commandline + +```bash +./build/tests/unittests/runUnitTests +``` + +#### Run profilerfeaturetests on the commandline + +```bash +./build/tests/featuretests/profiler/runFeatureTests +``` + +#### Run tracer featuretests on the commandline + +```bash +./build/tests/featuretests/tracer/runTracerFeatureTests +``` + +#### Run all tests + +```bash +rocprofv2 -t +``` + +OR + +```bash +ctest +``` + +### Guidelines for adding new tests + +- Prefer to enhance an existing test as opposed to writing a new one. Tests have overhead to start and many small tests spend precious test time on startup and initialization issues. +- Make the test run standalone without requirement for command-line arguments. This makes it easier to debug since the name of the test is shown in the test report and if you know the name of the test you can the run the test. + +## Logging + +To enable error messages logging to '/tmp/rocprofiler_log.txt': + +```bash +export ROCPROFILER_LOG=1 +``` + +## Kernel Name Truncation + +By default kernel names are not truncated. To enable truncation for readability: + +``` +export ROCPROFILER_TRUNCATE_KERNEL_PATH=1 +``` + +## Documentation + +We make use of doxygen to automatically generate API documentation. Generated document can be found in the following path: + +ROCM_PATH by default is /opt/rocm +It can be set by the user in different location if needed. +/share/doc/rocprofv2 + +installation: + +```string +rocprofiler-docs_9.0.0-local_amd64.deb +rocprofiler-docs-9.0.0-local.x86_64.rpm +``` + +## Samples + +- Profiling: Profiling Samples depending on replay mode +- Tracing: Tracing Samples + +installation: + +```string +rocprofiler-samples_9.0.0-local_amd64.deb +rocprofiler-samples-9.0.0-local.x86_64.rpm +``` + +usage: + +samples can be run as independent executables once installed + +## Project Structure + +- bin: ROCProf scripts along with V1 post processing scripts +- doc: Documentation settings for doxygen, V1 API Specifications pdf document. +- include: + - rocprofiler.h: V1 API Header File + - v2: + - rocprofiler.h: V2 API Header File + - rocprofiler_plugin.h: V2 Tool Plugins API +- plugin + - file: File Plugin + - perfetto: Perfetto Plugin + - ctf: CTF Plugin +- samples: Samples of how to use the API, and also input.txt input file samples for counter collection. +- script: Scripts needed for tracing +- src: Source files of the project + - api: API implementation for rocprofv2 + - core: Core source files needed for the V1/V2 API + - counters: Basic and Derived Counters + - hardware: Hardware support + - hsa: Provides support for profiler and tracer to communicate with HSA + - queues: Intercepting HSA Queues + - packets: Packets Preparation for profiling + - memory: Memory Pool used in buffers that saves the output data + - session: Session Logic + - filter: Type of profiling or tracing and its properties + - tracer: Tracing support of the session + - profiler: Profiling support of the session + - tools: Tools needed to run profiling and tracing + - rocsys: Controlling Session from another CLI + - utils: Utilities needed by the project +- tests: Tests folder +- CMakeLists.txt: Handles cmake list for the whole project +- build.sh: To easily build and compile rocprofiler +- CHANGELOG.md: Changes that are happening per release + +## Support + +Please report in the Github Issues. + +## Limitations + +- Navi3x requires a stable power state for counter collection. + Currently, this state needs to be set by the user. + To do so, set "power_dpm_force_performance_level" to be writeable for non-root users, then set performance level to profile_standard: + + ```bash + sudo chmod 777 /sys/class/drm/card0/device/power_dpm_force_performance_level + echo profile_standard >> /sys/class/drm/card0/device/power_dpm_force_performance_level + ``` + + Recommended: "profile_standard" for counter collection and "auto" for all other profiling. Use rocm-smi to verify the current power state. For multiGPU systems (includes integrated graphics), replace "card0" by the desired card. +- Timestamps may be incorrect with HIP_OPS when the system has been in sleep state. +- HIP_OPS are mutually exclusive with HSA_OPS. +- JSON Plugin as of now doesn't have an automated way to merge for multiple processes, a file will be generated per process (rank). + +## Supported AMD GPU Architectures (V2) + + The following AMD GPU architectures are supported with ROCprofiler V2: + +- gfx900 (AMD Vega 10) +- gfx906 (AMD Vega 7nm also referred to as AMD Vega 20) +- gfx908 (AMD Instinct™ MI100 accelerator) +- gfx90a (AMD Instinct™ MI200) +- gfx94x (AMD Instinct™ MI300) +- gfx10xx ([Navi2x] AMD Radeon(TM) Graphics) +- gfx11xx ([Navi3x] AMD Radeon(TM) Graphics) + +> [!WARNING] +> The latest mainline version of AQLprofile can be found at [https://repo.radeon.com/rocm/misc/aqlprofile/](https://repo.radeon.com/rocm/misc/aqlprofile/). However, it's important to note that updates to the public AQLProfile may not occur as frequently as updates to the rocprofiler. This discrepancy could lead to a potential mismatch between the AQLprofile binary and the rocprofiler source. diff --git a/projects/rocprofiler/RPM/post.in b/projects/rocprofiler/RPM/post.in new file mode 100644 index 0000000000..3fb4dd57af --- /dev/null +++ b/projects/rocprofiler/RPM/post.in @@ -0,0 +1,5 @@ +# left-hand term originates from ENABLE_LDCONFIG = ON/OFF at package build +if [ "@ENABLE_LDCONFIG@" == "ON" ]; then + echo @CPACK_PACKAGING_INSTALL_PREFIX@/@CMAKE_INSTALL_LIBDIR@ > /@CMAKE_INSTALL_SYSCONFDIR@/ld.so.conf.d/librocprofiler64.conf + ldconfig +fi diff --git a/projects/rocprofiler/RPM/postun.in b/projects/rocprofiler/RPM/postun.in new file mode 100644 index 0000000000..c5750b65e2 --- /dev/null +++ b/projects/rocprofiler/RPM/postun.in @@ -0,0 +1,4 @@ +# left-hand term originates from ENABLE_LDCONFIG = ON/OFF at package build +if [ "@ENABLE_LDCONFIG@" == "ON" ]; then + rm -f /@CMAKE_INSTALL_SYSCONFDIR@/ld.so.conf.d/librocprofiler64.conf && ldconfig +fi diff --git a/projects/rocprofiler/bin/build_kernel.sh b/projects/rocprofiler/bin/build_kernel.sh new file mode 100755 index 0000000000..318178156f --- /dev/null +++ b/projects/rocprofiler/bin/build_kernel.sh @@ -0,0 +1,84 @@ +#!/bin/sh -x +SO_EXT="hsaco" + +TEST_NAME=$1 +DST_DIR=$2 +ROCM_DIR=$3 +TGT_LIST=$4 + +if [ -z "$TEST_NAME" ] ; then + echo "Usage: $0 " + echo " Will look for .cl and will build .$SO_EXT dynamic code object library" + exit 1 +fi +OBJ_NAME=$(echo "_$(basename $TEST_NAME)" | sed -e 's/_./\U&\E/g' -e 's/_//g') + +if [ -z "$DST_DIR" ] ; then + DST_DIR=$(dirname TEST_NAME) +fi + +if [ -z "$ROCM_DIR" ] ; then + ROCM_DIR=/opt/rocm +fi + +if [ -z "$TGT_LIST" ] ; then + TGT_LIST=`$ROCM_DIR/bin/rocminfo | grep "amdgcn-amd-amdhsa--" | head -n 1 | sed -n "s/^.*amdgcn-amd-amdhsa--\(\w*\).*$/\1/p"` +fi + +if [ -z "$TGT_LIST" ] ; then + echo "Error: GPU targets not found" + exit 1 +fi + +OCL_VER="2.0" + +if [ -e $ROCM_DIR/llvm ] ; then + LLVM_DIR=$ROCM_DIR/llvm + LIB_DIR=$ROCM_DIR/lib +else + LLVM_DIR=$ROCM_DIR/hcc + LIB_DIR=$LLVM_DIR/lib +fi + +# Determine whether using new or old device-libs layout +if [ -e $LIB_DIR/bitcode/opencl.amdgcn.bc ]; then + BC_DIR=$LIB_DIR/bitcode +elif [ -e $LIB_DIR/opencl.amdgcn.bc ]; then + BC_DIR=$LIB_DIR +elif [ -e $ROCM_DIR/amdgcn/bitcode/opencl.bc ]; then + BC_DIR=$ROCM_DIR/amdgcn/bitcode +else + echo "Error: Cannot find amdgcn bitcode directory" + exit 1 +fi + +CLANG_ROOT=$LLVM_DIR/lib/clang +CLANG_DIR=`ls -d $CLANG_ROOT/* | head -n 1` +if [ "$CLANG_DIR" = "" ] ; then + echo "Error: LLVM clang library was not found" + exit 1 +fi + +BIN_DIR=$LLVM_DIR/bin +INC_DIR=$CLANG_DIR/include +if [ -e $BC_DIR/opencl.amdgcn.bc ]; then + BITCODE_OPTS="-nogpulib \ + -Xclang -mlink-bitcode-file -Xclang $BC_DIR/opencl.amdgcn.bc \ + -Xclang -mlink-bitcode-file -Xclang $BC_DIR/ockl.amdgcn.bc \ + -Xclang -mlink-bitcode-file -Xclang $BC_DIR/ocml.amdgcn.bc" +else + BITCODE_OPTS="--hip-device-lib-path=$BC_DIR" +fi + +for GFXIP in $TGT_LIST ; do + OBJ_PREF=$GFXIP + OBJ_FILE="${OBJ_PREF}_${OBJ_NAME}.$SO_EXT" + # Kernels are currently compiled with code object version 4. In order to + # switch to code object version 5, kernel argument setup in + # rocprofiler/test/simple_convolution/simple_convolution needs to be updated + # according to https://llvm.org/docs/AMDGPUUsage.html#amdgpu-amdhsa-code-object-kernel-argument-metadata-map-table-v5 + $BIN_DIR/clang -cl-std=CL$OCL_VER -include $INC_DIR/opencl-c.h $BITCODE_OPTS -target amdgcn-amd-amdhsa -mcpu=$GFXIP -mcode-object-version=4 $TEST_NAME.cl -o $DST_DIR/$OBJ_FILE + echo "'$OBJ_FILE' generated" +done + +exit 0 diff --git a/projects/rocprofiler/bin/dform.py b/projects/rocprofiler/bin/dform.py new file mode 100644 index 0000000000..97dbc62107 --- /dev/null +++ b/projects/rocprofiler/bin/dform.py @@ -0,0 +1,107 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +import os +from sqlitedb import SQLiteDB + + +def gen_message(outfile): + if outfile != "": + print("File '" + outfile + "' is generating") + + +def post_process_data(db, table_name, outfile=""): + # db.add_data_column('A', 'DispDurNs', 'INTEGER', 'BeginNs - DispatchNs') + # db.add_data_column('A', 'ComplDurNs', 'INTEGER', 'CompleteNs - EndNs') + # db.add_data_column('A', 'TotalDurNs', 'INTEGER', 'CompleteNs - DispatchNs') + # db.add_data_column(table_name, 'TimeNs', 'INTEGER', 'BeginNs - %d' % start_ns) + db.add_data_column(table_name, "DurationNs", "INTEGER", "EndNs - BeginNs") + if outfile != "": + db.dump_csv(table_name, outfile) + gen_message(outfile) + + +def gen_data_bins(db, outfile): + db.execute( + "create view C as select Name, Calls, TotalDurationNs, TotalDurationNs/Calls as AverageNs, TotalDurationNs*100.0/(select sum(TotalDurationNs) from %s) as Percentage from %s order by TotalDurationNs desc;" + % ("B", "B") + ) + db.dump_csv("C", outfile) + db.execute("DROP VIEW C") + + +def gen_table_bins(db, table, outfile, name_var, dur_ns_var): + db.execute( + "create view B as select (%s) as Name, count(%s) as Calls, sum(%s) as TotalDurationNs from %s group by %s" + % (name_var, name_var, dur_ns_var, table, name_var) + ) + gen_data_bins(db, outfile) + db.execute("DROP VIEW B") + gen_message(outfile) + + +def gen_api_json_trace(db, table, start_ns, outfile): + db.execute( + 'create view B as select "Index", Name as name, __section as pid, __lane as tid, ((BeginNs - %d)/1000) as ts, (DurationNs/1000) as dur from %s;' + % (start_ns, table) + ) + db.dump_json("B", table, outfile) + db.execute("DROP VIEW B") + gen_message(outfile) + + +def gen_ext_json_trace(db, table, start_ns, outfile): + db.execute( + "create view B as select Name as name, __section as pid, __lane as tid, ((BeginNs - %d)/1000) as ts, ((EndNs - BeginNs)/1000) as dur from %s;" + % (start_ns, table) + ) + db.dump_json("B", table, outfile) + db.execute("DROP VIEW B") + gen_message(outfile) + + +def gen_ops_json_trace(db, table, base_pid, start_ns, outfile): + db.execute( + 'create view B as select "Index", "%s" as name, ("dev-id" + %d) as pid, __lane as tid, ((BeginNs - %d)/1000) as ts, (DurationNs/1000) as dur from %s;' + % ( + "roctx-range" if "ROCP_RENAME_KERNEL" in os.environ else "Name", + base_pid, + start_ns, + table, + ) + ) + db.dump_json("B", table, outfile) + db.execute("DROP VIEW B") + gen_message(outfile) + + +def gen_kernel_json_trace(db, table, base_pid, start_ns, outfile): + db.execute( + 'create view B as select "Index", KernelName as name, ("gpu-id" + %d) as pid, tid, ((BeginNs - %d)/1000) as ts, (DurationNs/1000) as dur from %s;' + % (base_pid, start_ns, table) + ) + db.dump_json("B", table, outfile) + db.execute("DROP VIEW B") + gen_message(outfile) + + +############################################################################################## diff --git a/projects/rocprofiler/bin/mem_manager.py b/projects/rocprofiler/bin/mem_manager.py new file mode 100755 index 0000000000..cc04d2343b --- /dev/null +++ b/projects/rocprofiler/bin/mem_manager.py @@ -0,0 +1,485 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +import sys, os, re +from sqlitedb import SQLiteDB + +pinned = ["hipMallocHost", "hipHostMalloc", "hipHostAlloc"] +ondevice = [ + "hipMalloc", + "hipMallocPitch", + "hipMallocArray", + "hipMalloc3DArray", + "hsa_amd_memory_pool_allocate", +] + +mm_table_descr = [ + [ + "BeginNs", + "EndNs", + "pid", + "tid", + "Name", + "Direction", + "SrcType", + "DstType", + "Size", + "Async", + ], + { + "BeginNs": "INTEGER", + "EndNs": "INTEGER", + "pid": "INTEGER", + "tid": "INTEGER", + "Name": "TEXT", + "Direction": "TEXT", + "SrcType": "TEXT", + "DstType": "TEXT", + "Size": "INTEGER", + "Async": "TEXT", + }, +] + + +def fatal(msg): + sys.stderr.write(sys.argv[0] + ": " + msg + "\n") + sys.exit(1) + + +DELIM = "," + + +# Mem copy manager class +class MemManager: + def __init__(self, db, indir): + self.db = db + self.allocations = {} + self.hsa_agent_types = {} + self.memcopies = {} + self.filename = "" + self.fd = 0 + self.parse_hsa_handles(indir + "/" + "hsa_handles.txt") + + def __del__(self): + if self.fd != 0: + self.fd.close() + + # Parsing the mapping of HSA agent and memory pool handles + def parse_hsa_handles(self, infile): + if os.path.exists(infile): + inp = open(infile, "r") + cpu_agent_ptrn = re.compile(r"(0x[0-9a-fA-F]+) agent cpu") + gpu_agent_ptrn = re.compile(r"(0x[0-9a-fA-F]+) agent gpu") + for line in inp.readlines(): + m_cpu = cpu_agent_ptrn.match(line) + if m_cpu: + self.hsa_agent_types[str(int(m_cpu.group(1), 16))] = 0 # "cpu" + m_gpu = gpu_agent_ptrn.match(line) + if m_gpu: + self.hsa_agent_types[str(int(m_gpu.group(1), 16))] = 1 # "gpu" + inp.close() + + # register alloc and memcpy API calls + # ['BeginNs', 'EndNs', 'pid', 'tid', 'Name', 'args', 'Index', 'Data'], + def register_api(self, rec_vals): + res = "" + record_name = rec_vals[4] # 'Name' + record_args = rec_vals[5] # 'args' + malloc_ptrn = re.compile(r"hip.*Malloc|hsa_amd_memory_pool_allocate") + mcopy_ptrn = re.compile(r"hipMemcpy|hsa_amd_memory_async_copy") + + if malloc_ptrn.match(record_name): + self.add_allocation(record_name, record_args) + elif mcopy_ptrn.match(record_name): + res = self.add_memcpy(rec_vals) + + return res + + # register memcpy asynchronous copy + # ['BeginNs', 'EndNs', 'Name', 'pid', 'tid', 'Index', ... + def register_copy(self, rec_vals): + data = "" + event = rec_vals[2] # 'Name' + procid = rec_vals[3] # 'pid' + recordid = rec_vals[5] # 'Index' + # query syncronous memcopy API record + key = (recordid, procid, 0) + if key in self.memcopies: + data = self.memcopies[key] + + # query asyncronous memcopy API record + key = (recordid, procid, 1) + if key in self.memcopies: + if data != "": + fatal("register_copy: corrupted record sync/async") + async_copy_start_time = rec_vals[0] + async_copy_end_time = rec_vals[1] + + tid = rec_vals[4] + copy_line_header = ( + str(async_copy_start_time) + + DELIM + + str(async_copy_end_time) + + DELIM + + str(procid) + + DELIM + + str(tid) + ) + copy_line_footer = "Async=" + str(1) + data = copy_line_header + self.memcopies[key] + copy_line_footer + self.memcopies[key] = data + + return data + + # register memcpy asynchronous activity + # rec_vals: ['BeginNs', 'EndNs', 'dev-id', 'queue-id', 'Name', 'pid', 'tid', 'Index', 'Data', ... + def register_activity(self, rec_vals): + data = "" + procid = rec_vals[5] # 'pid' + recordid = rec_vals[7] # 'Index' + + # query syncronous memcopy API record + key = (recordid, procid, 0) + if key in self.memcopies: + data = self.memcopies[key] + + # query asyncronous memcopy API record + key = (recordid, procid, 1) + if key in self.memcopies: + if data != "": + fatal("register_activity: corrupted record sync/async") + + async_copy_start_time = rec_vals[0] + async_copy_end_time = rec_vals[1] + + tid = rec_vals[6] + copy_line_header = ( + str(async_copy_start_time) + + DELIM + + str(async_copy_end_time) + + DELIM + + str(procid) + + DELIM + + str(tid) + ) + copy_line_footer = "Async=" + str(1) + data = copy_line_header + self.memcopies[key] + copy_line_footer + self.memcopies[key] = data + + return data + + # add allocation to map + def add_allocation(self, event, args): + choice = 0 + if event == "hipMallocPitch": + malloc_args_ptrn = re.compile(r"\(ptr\((.*)\) width\((.*)\) height\((.*)\)\)") + choice = 1 + elif event == "hipMallocArray": + malloc_args_ptrn = re.compile( + r"\(array\((.*)\) width\((.*)\) height\((.*)\)\)" + ) + choice = 1 + elif event == "hipMalloc3DArray": + malloc_args_ptrn = re.compile( + r"\(array\((.*)\) width\((.*)\) height\((.*)\) depth\((.*)\)\)" + ) + choice = 2 + elif event == "hsa_amd_memory_pool_allocate": + # ({handle=25291264}, 40, 0, 0x7ffc4c7bf1b0) + malloc_args_ptrn = re.compile( + r"\({handle=\d+}, (\d+), \d+, (0x[0-9a-fA-F]+)\)" + ) + choice = 4 + else: + # (ptr(0x7f3407000000) size(800000000) flags(0)) + malloc_args_ptrn = re.compile(r"\(ptr\((.*)\) size\((.*)\) .*\)") + choice = 3 + m = malloc_args_ptrn.match(args) + if m: + if choice == 4: + ptr = int(m.group(2), 16) + size = int(m.group(1)) + elif choice == 3: + ptr = int(m.group(1), 16) + size = int(m.group(2)) + elif choice == 1: + ptr = int(m.group(1), 16) + size = int(m.group(2)) * int(m.group(3)) + else: + ptr = int(m.group(1), 16) + size = int(m.group(2)) * int(m.group(3)) * int(m.group(4)) + self.allocations[ptr] = (size, event) + + # get type of ptr + def get_ptr_type(self, ptr): + addr = int(ptr, 16) + addr_type = "unknown" + found = 0 + for base, (size, event) in self.allocations.items(): + if addr >= base and addr < base + size: + found = 1 + break + if not found: + addr_type = "pageable" + elif event in pinned: + addr_type = "pinned" + elif event in ondevice: + addr_type = "device" + elif ptr in self.hsa_agent_types: + if self.hsa_agent_types[ptr] == 0: + addr_type = "pinned" + elif self.hsa_agent_types[ptr] == 1: + addr_type = "device" + else: + fatal("internal error: ptr(" + ptr + ") cannot be identified") + else: + fatal("internal error: ptr(" + ptr + ") cannot be identified") + return addr_type + + # add memcpy to map + def add_memcpy(self, recvals): + recordid = recvals[6] # same as corrid + event = recvals[4] + start_time = recvals[0] # sync time stamp + end_time = recvals[1] # sync time stamp + args = recvals[5] + procid = int(recvals[2]) # used to query async entries + pid = recvals[2] + tid = recvals[3] + + # hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) + hip_memcpy_ptrn = re.compile( + r"\(\s*dst\((.*)\) src\((.*)\) sizeBytes\((\d+)\).*\)" + ) + # hipMemcpy2D(void* dst, size_t dpitch, const void* src, size_t spitch, size_t width, + # size_t height, hipMemcpyKind kind); + hip_memcpy_ptrn2 = re.compile( + r"\(\s*dst\((.*)\) .* src\((.*)\) .* width\((\d+)\) height\((\d+)\).*\)" + ) + # hipMemcpyToArray(hipArray* dst, size_t wOffset, size_t hOffset, const void* src, + # size_t count, hipMemcpyKind kind); + hip_memcpy_ptrn3 = re.compile( + r"\(\s*dst\((.*)\) .* src\((.*)\) count\((\d+)\).*\)" + ) + # hipMemcpyToSymbol(const void* symbolName, const void* src, size_t sizeBytes, + # size_t offset = 0, hipMemcpyKind kind) + hip_memcpy_ptrn4 = re.compile( + r"\(\s*symbol\((.*)\) src\((.*)\) sizeBytes\((\d+)\).*\)" + ) + # memcopy with kind argument + hip_memcpy_ptrn_kind = re.compile(r".* kind\((\d+)\)\s*.*") + # hsa_amd_memory_async_copy(void* dst, hsa_agent_t dst_agent, const void* src, + # hsa_agent_t src_agent, size_t size, + # uint32_t num_dep_signals, + # const hsa_signal_t* dep_signals, + # hsa_signal_t completion_signal); + # "(0x7f8ab6600000, 27064880, 0x7f8b16000000, 27059968, 800000000, 0, 0, 140240759809536) = 0" + # hsa_memcpy_ptrn_prev used to support format transition and will be cleaned up later. + hsa_memcpy_ptrn_prev = re.compile( + r"\((0x[0-9a-fA-F]+), (\d+), (0x[0-9a-fA-F]+), (\d+), (\d+), .*\) = \d" + ) + # "(0x7fd83bc00000, {handle=16124864}, 0x7fd89b600000, {handle=16119808}, 800000000, 0, 0, {handle=140573877724672}) = 0" + hsa_memcpy_ptrn = re.compile( + r"\((0x[0-9a-fA-F]+), {handle=(\d+)}, (0x[0-9a-fA-F]+), {handle=(\d+)}, (\d+), .*\) = \d" + ) + # "(0x7f9125cfe7b0, 0x7f9125cfe784, 0x7f9125cfe790, 0x7f9125cfe784, 0x7f9125cfe778, {handle=94324038652880}, 1, 0, 0, {handle=140261380710784}) = 0" + # dst, dst_offset, src, src_offset, range, copy_agent, dir, num_dep_signals, dep_signals, completion_signal + hsa_memcpy_ptrn2 = re.compile( + r"\((0x[0-9a-fA-F]+), 0x[0-9a-fA-F]+, (0x[0-9a-fA-F]+), 0x[0-9a-fA-F]+, 0x[0-9a-fA-F]+, {z=(\d+), y=(\d+), x=(\d+)}, {handle=(\d+)}, .*\) = \d" + ) + # aysnc memcopy + async_event_ptrn = re.compile(r"Async|async") + m_basic_hip = hip_memcpy_ptrn.match(args) + m_basic_hsa3 = hip_memcpy_ptrn4.match(args) + m_basic_hsa_prev = hsa_memcpy_ptrn_prev.match(args) + m_basic_hsa = hsa_memcpy_ptrn.match(args) + m_basic_hsa2 = hsa_memcpy_ptrn2.match(args) + is_hip = True if not (m_basic_hsa_prev or m_basic_hsa or m_basic_hsa2) else False + m_2d = hip_memcpy_ptrn2.match(args) + m_array = hip_memcpy_ptrn3.match(args) + is_async = 1 if async_event_ptrn.search(event) else 0 + copy_line = "" + size = 0 + dstptr_type = "unknown" + srcptr_type = "unknown" + direction = "unknown" + kind_switcher = { + "0": "HtoH", + "1": "HtoD", + "2": "DtoH", + "3": "DtoD", + "4": "auto", + } + + condition_matched = False + if m_basic_hip: + dstptr = m_basic_hip.group(1) + dstptr_type = self.get_ptr_type(dstptr) + srcptr = m_basic_hip.group(2) + srcptr_type = self.get_ptr_type(srcptr) + size = int(m_basic_hip.group(3)) + condition_matched = True + + if m_basic_hsa_prev: + dstptr = m_basic_hsa_prev.group(1) + dst_agent_ptr = m_basic_hsa_prev.group(2) + dstptr_type = self.get_ptr_type(dst_agent_ptr) + srcptr = m_basic_hsa_prev.group(3) + src_agent_ptr = m_basic_hsa_prev.group(4) + srcptr_type = self.get_ptr_type(src_agent_ptr) + size = int(m_basic_hsa_prev.group(5)) + condition_matched = True + + if m_basic_hsa: + dstptr = m_basic_hsa.group(1) + dst_agent_ptr = m_basic_hsa.group(2) + dstptr_type = self.get_ptr_type(dst_agent_ptr) + srcptr = m_basic_hsa.group(3) + src_agent_ptr = m_basic_hsa.group(4) + srcptr_type = self.get_ptr_type(src_agent_ptr) + size = int(m_basic_hsa.group(5)) + condition_matched = True + + if m_basic_hsa2: + dstptr = m_basic_hsa2.group(1) + dst_agent_ptr = m_basic_hsa2.group(6) + dstptr_type = self.get_ptr_type(dst_agent_ptr) + srcptr = m_basic_hsa2.group(2) + src_agent_ptr = m_basic_hsa2.group(6) + srcptr_type = self.get_ptr_type(src_agent_ptr) + z = int(m_basic_hsa2.group(3)) + y = int(m_basic_hsa2.group(4)) + x = int(m_basic_hsa2.group(5)) + size = x * y * z + condition_matched = True + + if m_basic_hsa3: + dstptr = m_basic_hsa3.group(1) + dstptr_type = self.get_ptr_type(dstptr) + srcptr = m_basic_hsa3.group(2) + srcptr_type = self.get_ptr_type(srcptr) + size = int(m_basic_hsa3.group(3)) + condition_matched = True + + if m_array: + dstptr = m_array.group(1) + dstptr_type = self.get_ptr_type(dstptr) + srcptr = m_array.group(2) + srcptr_type = self.get_ptr_type(srcptr) + size = m_array.group(3) + condition_matched = True + + if m_2d: + dstptr = m_2d.group(1) + dstptr_type = self.get_ptr_type(dstptr) + srcptr = m_2d.group(2) + srcptr_type = self.get_ptr_type(srcptr) + size = int(m_2d.group(3)) * int(m_2d.group(4)) + condition_matched = True + + if not condition_matched: + fatal('Memcpy args "' + args + '" cannot be identified') + + if not is_async: + start_time = recvals[0] # sync time stamp + end_time = recvals[1] # sync time stamp + duration = int(end_time) - int(start_time) + + evt_switcher = { + "hipMemcpyDtoD": "DtoD", + "hipMemcpyDtoDAsync": "DtoD", + "hipMemcpyDtoH": "DtoH", + "hipMemcpyDtoHAsync": "DtoH", + "hipMemcpyHtoD": "HtoD", + "hipMemcpyHtoDAsync": "HtoD", + } + + if is_hip: + m = hip_memcpy_ptrn_kind.match(args) + if m: + direction = kind_switcher.get(m.group(1), "unknown") + else: + direction = evt_switcher.get(event, "unknown") + else: + if ( + dst_agent_ptr in self.hsa_agent_types + and src_agent_ptr in self.hsa_agent_types + ): + if self.hsa_agent_types[src_agent_ptr] == 1: + direction = "D" + elif self.hsa_agent_types[src_agent_ptr] == 0: + direction = "H" + if direction != "unknown": + direction += "to" + if self.hsa_agent_types[dst_agent_ptr] == 1: + direction += "D" + elif self.hsa_agent_types[dst_agent_ptr] == 0: + direction += "H" + + copy_line_header = "" + copy_line_footer = "" + copy_line_header = ( + str(start_time) + DELIM + str(end_time) + DELIM + str(pid) + DELIM + str(tid) + ) + copy_line_footer = "Async=" + str(is_async) + + copy_line = ( + copy_line_header + + DELIM + + event + + DELIM + + "Direction=" + + direction + + DELIM + + "SrcType=" + + srcptr_type + + DELIM + + "DstType=" + + dstptr_type + + DELIM + + "Size=" + + str(size) + + DELIM + + copy_line_footer + ) + + self.memcopies[(recordid, procid, is_async)] = copy_line + return copy_line + + def dump_data(self, table_name, file_name): + # To create memcopy info table in DB + print("File '" + file_name + "' is generating") + table_handle = self.db.add_table(table_name, mm_table_descr) + + fld_ptrn = re.compile(r"(.*)=(.*)") + for key, record in self.memcopies.items(): + rec_vals_array = [] + for rec in record.split(DELIM): + fld_ptrnm = fld_ptrn.match(rec) + if fld_ptrnm: + rec_vals_array.append(fld_ptrnm.group(2)) + else: + rec_vals_array.append(rec) + self.db.insert_entry(table_handle, rec_vals_array) + + # To dump the memcopy info table as CSV + self.db.dump_csv(table_name, file_name) diff --git a/projects/rocprofiler/bin/merge_traces.sh b/projects/rocprofiler/bin/merge_traces.sh new file mode 100755 index 0000000000..1c6e65bf6b --- /dev/null +++ b/projects/rocprofiler/bin/merge_traces.sh @@ -0,0 +1,102 @@ +#!/bin/bash + +################################################################################ +# Copyright (c) 2021 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. +################################################################################ + +#A script to merge rocprof traces and then provide a results.json for the aggregate numbers. + +BIN_DIR=$(dirname $(realpath ${BASH_SOURCE[0]})) + +bin_name=`basename $0` + +# usage method +usage() { + echo "Script for aggregating results from multiple rocprofiler out directries." + echo "Full path: $BIN_DIR/$bin_name" + echo "Usage: if running independently" + echo " $bin_name -o [...]" + echo "" + echo "Usage: if running with rocprof" + echo " rocprof --merge-traces -o [...]" + echo "" + echo "Options:" + echo " -o - output directory where the results will be aggregated." + echo " ... - space separated list of rocprofiler directories. If not specified, CWD is used." + echo "" + exit 1 +} + +# read arguments + +INPUT_DIRS=() +while getopts "o:h" opt; do + case $opt in + o) OUTPUT_DIR=$OPTARG ;; + h) usage ;; + \?) usage ;; +esac +done +shift $((OPTIND-1)) + +INPUT_DIRS=$@ + +if [ "${OUTPUT_DIR}" = "" ] ; then + echo "Missing output dir option" + usage +fi + +for INPUT_DIR in ${INPUT_DIRS} ; do + if [[ ! -d "${INPUT_DIR}" ]] ; then + echo "Directory ${INPUT_DIR} does not exist." + exit 1 + fi +done + +if ! [ -d "${OUTPUT_DIR}" ] ; then + mkdir -p "${OUTPUT_DIR}" +fi + +echo "Processing directories: $INPUT_DIRS" +for file in begin_ts_file hcc_ops_trace hsa_handles hip_api_trace roctx_trace hsa_api_trace results async_copy_trace; do + res=$(find ${INPUT_DIRS} -type f -regextype sed -regex ".*/[0-9]\{1,\}_${file}\.txt" \ + -not -path "${OUTPUT_DIR}/*") + test -n "${res}" && cat ${res} > "${OUTPUT_DIR}/${file}.txt" +done + +if ! [ -d "${BIN_DIR}" ] ; then + echo "Bin directory $BIN_DIR not found!" + exit 1 +fi + +if [ -z "$ROCP_PYTHON_VERSION" ] ; then + ROCP_PYTHON_VERSION=python3 +fi + +OUTPUT_LIST="$OUTPUT_DIR/results.txt" +db_output="$OUTPUT_DIR/results.db" +echo "$ROCP_PYTHON_VERSION $BIN_DIR/tblextr.py $db_output $OUTPUT_LIST" +$ROCP_PYTHON_VERSION $BIN_DIR/tblextr.py $db_output $OUTPUT_LIST +if [ "$?" -ne 0 ] ; then + echo "Profiling data corrupted: '$OUTPUT_LIST'" + exit 1 +fi + diff --git a/projects/rocprofiler/bin/rocprofv2 b/projects/rocprofiler/bin/rocprofv2 new file mode 100755 index 0000000000..95c72bc092 --- /dev/null +++ b/projects/rocprofiler/bin/rocprofv2 @@ -0,0 +1,351 @@ +#!/bin/bash -e + +set -eo pipefail + +# Define color code +GREEN='\033[0;32m' +GREY='\033[0;90m' +RED='\033[0;31m' +RESET='\033[0m' + +# Function to display deprecation warning message +display_warning() { + echo -e "\n${RED}WARNING: We are phasing out development and support for roctracer/rocprofiler/rocprof/rocprofv2 in favor of \ +rocprofiler-sdk/rocprofv3 in upcoming ROCm releases. Going forward, only critical defect fixes will be addressed for \ +older versions of profiling tools and libraries. We encourage all users to upgrade to the latest version, \ +rocprofiler-sdk library and rocprofv3 tool, to ensure continued support and access to new features.${RESET}\n" +} + +# Display the warning message +display_warning + +# LD_PRELOAD on script will not get propagated +if [ -n "${ROCP_PRELOAD}" ]; then LD_PRELOAD="${ROCP_PRELOAD}"; fi + +CURRENT_DIR="$( dirname -- "$0"; )"; +ROCPROFV2_DIR=$(dirname -- $(realpath ${BASH_SOURCE[0]})); +ROCM_DIR=$( dirname -- "$ROCPROFV2_DIR"; ) +PLUGIN_LIST=("ctf" "perfetto" "file" "att" "cli" "json") +RUN_FROM_BUILD=0 +if [[ $ROCPROFV2_DIR == *"/build"* ]]; then + RUN_FROM_BUILD=1 + ROCM_DIR=$ROCPROFV2_DIR +fi + +export ROCPROFILER_METRICS_PATH=$ROCM_DIR/libexec/rocprofiler/counters/derived_counters.xml +export LD_LIBRARY_PATH=$ROCM_DIR/lib:$LD_LIBRARY_PATH +export HSA_TOOLS_LIB=$ROCM_DIR/lib/librocprofiler64.so.2 + + +usage() { + echo -e "${RESET}ROCProfilerV2 Run Script Usage:" + echo -e "${GREEN}-h | --help ${RESET} For showing this message" + echo -e "${GREEN}--list-counters ${RESET} For showing all available counters for the current GPUs" + if [ $RUN_FROM_BUILD == 1 ]; then + echo -e "${GREEN}-t | --test ${RESET} For Running the tests" + echo -e "${GREEN}-mt | --mem-test ${RESET} For Running the Memory Leak tests. This run requires building using -acb | --asan-clean-build option" + fi + echo -e "${GREEN}-m ${RESET} For providing an absolute path of a custom metrics file" + echo -e "${GREEN}--basenames ${RESET} For Truncating the kernel names" + echo -e "${GREEN}--hip-api ${RESET} For Collecting HIP API Traces" + echo -e "${GREEN}--hip-activity | --hip-trace ${RESET} For Collecting HIP API Activities Traces" + echo -e "${GREEN}--hsa-api ${RESET} For Collecting HSA API Traces" + echo -e "${GREEN}--hsa-activity | --hsa-trace ${RESET} For Collecting HSA API Activities Traces" + echo -e "${GREEN}--roctx-trace ${RESET} For Collecting ROCTx Traces" + echo -e "${GREEN}--kernel-trace ${RESET} For Collecting Kernel dispatch Traces" + echo -e "${GREEN}--sys-trace ${RESET} For Collecting HIP and HSA APIs and their Activities Traces along ROCTX\n" + echo -e "\t#${GREY}usage e.g: rocprofv2 --[hip-trace|hsa-trace|roctx-trace|kernel-trace|sys-trace] \n"${RESET} + echo -e "${GREEN}--plugin ${RESET} PLUGIN_NAME For enabling a plugin (cli/file/perfetto/ctf/speedscope)" + echo -e "\t#${GREY} usage(file/perfetto/ctf) e.g: rocprofv2 -i pmc.txt --plugin [file/perfetto/ctf/json] -d out_dir " + echo -e "\t# use \"rocprofv2 --plugin json --disable-json-data-flows ...\" for SpeedScope support as speedscope doesn't support data flows.${RESET}\n" + echo -e "${GREEN}--plugin-version ${RESET} <1|2> For selecting the version for the plugin (1/2)" + echo -e "\t#${GREY} 1 - Legacy output format, 2 - New output format (default)${RESET}\n" + echo -e "${GREEN}-i | --input ${RESET} For adding counters file path (every line in the text file represents a counter)" + echo -e "\t#${GREY} usage: rocprofv2 -i pmc.txt -d ${RESET}\n" + echo -e "${GREEN}-o | --output-file-name ${RESET} For the output file name" + echo -e "\t#${GREY} usage e.g:(with current dir): rocprofv2 --hip-trace -o " + echo -e "\t#${GREY} usage e.g:(with custom dir): rocprofv2 --hip-trace -d -o ${RESET}\n" + echo -e "${GREEN}-d | --output-directory ${RESET} For adding output path where the output files will be saved" + echo -e "\t#${GREY} usage e.g:(with custom dir): rocprofv2 --hip-trace -d ${RESET}\n" + echo -e "${GREEN}-fi | --flush-interval ${RESET} For adding a flush interval in milliseconds, every \"flush interval\" the buffers will be flushed" + echo -e "\t#${GREY} usage e.g: rocprofv2 --hip-trace -fi 1000 ${RESET}\n" + echo -e "${GREEN}-tp | --trace-period ${RESET} Specifies a trace period in milliseconds, with format \"-tp ::\"." + echo -e "\t#${GREY} usage e.g: rocprofv2 --hip-trace -tp 1000:2000:4000 ${RESET}\n" + echo -e "${GREEN}-ns | --no-serialization ${RESET} For disabling serilization when running in counter-collection mode\"." + echo -e "\t#${GREY} usage e.g: rocprofv2 -i pmc.txt -ns${RESET}\n" + exit 1 +} + +if [ -z "$1" ]; then + usage + exit 1 +fi + +OUTPUT_PATH_INTERNAL="." + +while [ 1 ]; do + if [[ "$1" == "-h" || "$1" == "--help" ]]; then + usage + exit 1 + elif [[ "$1" == "-t" || "$1" == "--test" ]]; then + if [ $RUN_FROM_BUILD == 1 ]; then + ./run_tests.sh + exit 1 + fi + elif [[ "$1" == "-mt" || "$1" == "--mem-test" ]]; then + if [ $RUN_FROM_BUILD == 1 ]; then + $ROCM_DIR/tests-v2/memorytests/run_asan_tests.sh $ROCM_DIR/tests-v2/featuretests/profiler/apps/hip_vectoradd $ROCM_DIR/memleaks.log + exit 1 + fi + elif [[ "$1" == "--list-counters" ]]; then + export LD_PRELOAD=$LD_PRELOAD:$ROCM_DIR/lib/rocprofiler/librocprofiler_tool.so + eval $ROCM_DIR/libexec/rocprofiler/ctrl + exit 1 + elif [[ "$1" == "-i" || "$1" == "--input" ]]; then + if [ $2 ] && [ -n $2 ] && [ -r $2 ]; then + export COUNTERS_PATH=$2 + else + echo -e "Error: \"$2\" doesn't exist!" + usage + exit 1 + fi + shift + shift + elif [[ "$1" == "-o" || "$1" == "--output-file-name" ]]; then + if [ $2 ]; then + export OUT_FILE_NAME=$2 + else + usage + exit 1 + fi + shift + shift + elif [[ "$1" == "-d" || "$1" == "--output-directory" ]]; then + if [ $2 ]; then + OUTPUT_PATH_INTERNAL=$2 + export OUTPUT_PATH=$OUTPUT_PATH_INTERNAL + else + usage + exit 1 + fi + shift + shift + elif [[ "$1" == "-fi" || "$1" == "--flush-interval" ]]; then + if [ $2 ] && [ $2 -gt 0 ]; then + export ROCPROFILER_FLUSH_INTERVAL=$2 + else + echo -e "Wrong input \"$2\" for flush interval, it needs to be integer greater than zero!" + usage + exit 1 + fi + shift + shift + elif [[ "$1" == "-tp" || "$1" == "--trace-period" ]]; then + if [ $2 ] && [[ "$2" == *":"* ]]; then + export ROCPROFILER_TRACE_PERIOD=$2 + else + echo -e "Wrong input \"$2\" for trace period!" + usage + exit 1 + fi + shift + shift + elif [ "$1" == "--hip-api" ]; then + export ROCPROFILER_HIP_API_TRACE=1 + shift + elif [[ "$1" == "-ns" || "$1" == "--no-serialization" ]]; then + export ROCPROFILER_NO_SERIALIZATION=1 + shift + elif [[ "$1" == "--hip-activity" || "$1" == "--hip-trace" ]]; then + export ROCPROFILER_HIP_API_TRACE=1 + export ROCPROFILER_HIP_ACTIVITY_TRACE=1 + shift + elif [ "$1" == "--hsa-api" ]; then + export ROCPROFILER_HSA_API_TRACE=1 + shift + elif [[ "$1" == "--hsa-activity" || "$1" == "--hsa-trace" ]]; then + export ROCPROFILER_HSA_API_TRACE=1 + export ROCPROFILER_HSA_ACTIVITY_TRACE=1 + shift + elif [ "$1" == "--roctx-trace" ]; then + export ROCPROFILER_ROCTX_TRACE=1 + shift + elif [ "$1" == "--kernel-trace" ]; then + export ROCPROFILER_KERNEL_TRACE=1 + shift + elif [ "$1" == "--sys-trace" ]; then + export ROCPROFILER_HIP_API_TRACE=1 + export ROCPROFILER_HIP_ACTIVITY_TRACE=1 + export ROCPROFILER_HSA_API_TRACE=1 + export ROCPROFILER_ROCTX_TRACE=1 + shift + elif [ "$1" == "--roc-sys" ]; then + export ROCPROFILER_ENABLE_ROCSYS=$2 + shift + shift + elif [ "$1" == "-m" ]; then + if [ $2 ] && [ -n $2 ] && [ -r $2 ]; then + export ROCPROFILER_METRICS_PATH=$2 + else + echo -e "Error: \"$2\" path doesn't exist!" + usage + exit 1 + fi + shift + shift + elif [ "$1" == "--version" ]; then + if [ -f "$ROCM_DIR/libexec/rocprofiler/rocprofiler-version" ]; then + ROCPROFILER_LIBRARY_VERSION=2 $ROCM_DIR/libexec/rocprofiler/rocprofiler-version + else + ROCM_VERSION=$(cat $ROCM_DIR/.info/version) + echo -e "ROCm version: $ROCM_VERSION" + echo -e "ROCProfiler version: 2.0" + fi + exit 0 + elif [ "$1" == "--basenames" ]; then + export ROCPROFILER_TRUNCATE_KERNEL_PATH=1 + shift + elif [ "$1" == "--tool-version" ]; then + shift + shift + elif [ "$1" == "--plugin-version" ]; then + export ROCPROFILER_PLUGIN_LIB_VERSION=".$2" + if [ ! -z $ROCPROFILER_PLUGIN_LIB ]; then + if [ -f $ROCM_DIR/lib/rocprofiler/$ROCPROFILER_PLUGIN_LIB$ROCPROFILER_PLUGIN_LIB_VERSION ]; then + export ROCPROFILER_PLUGIN_LIB=$ROCPROFILER_PLUGIN_LIB$ROCPROFILER_PLUGIN_LIB_VERSION + else + echo -e "Error: \"$ROCPROFILER_PLUGIN_LIB$ROCPROFILER_PLUGIN_LIB_VERSION\" doesn't exist!" + exit 1 + fi + fi + shift + shift + elif [ "$1" == "--plugin" ]; then + if [ -n $2 ]; then + PLUGIN=$2 + if [[ ! "${PLUGIN_LIST[*]}" =~ $PLUGIN ]]; then + echo -e "Wrong input \"$2\" for plugin!" + usage + exit 1 + fi + export ROCPROFILER_PLUGIN_LIB=lib${PLUGIN}_plugin.so$ROCPROFILER_PLUGIN_LIB_VERSION + else + echo -e "Wrong input \"$2\" for plugin!" + usage + exit 1 + fi + if [ "$2" = "att" ]; then + export ROCPROFILER_TRUNCATE_KERNEL_PATH=1 + if [ $RUN_FROM_BUILD == 1 ]; then + ATT_PATH=$ROCM_DIR/plugin/att/att/att.py + else + ATT_PATH=$ROCM_DIR/libexec/rocprofiler/att/att.py + export ROCPROFV2_ATT_LIB_PATH=$ROCM_DIR/lib/hsa-amd-aqlprofile/librocprofv2_att.so + fi + ATT_ARGV=$3 + shift + + ATT_OPTIONS="Not done" + ATT_PYTHON3_ARG="python3 " + while [ "$ATT_OPTIONS" = "Not done" ]; do + if [[ "$3" = "--trace_file" ]]; then + ATT_ARGV="$ATT_ARGV $3 \"$4\"" + shift + shift + elif [[ "$3" = "--mode" || "$3" = "--ports" || "$3" == "--att_kernel" ]]; then + ATT_ARGV="$ATT_ARGV $3 $4" + shift + shift + else + ATT_OPTIONS="Done" + fi + done + fi + shift + shift + elif [ "$1" == "--disable-json-data-flows" ]; then + export ROCPROFILER_DISABLE_JSON_DATA_FLOWS=1 + shift + elif [[ "$1" == "-"* || "$1" == "--"* ]]; then + echo -e "Wrong option \"$1\", Please use the following options:\n" + usage + exit 1 + else + break + fi +done + +PMC_LINES=() +if [ -n "$COUNTERS_PATH" ]; then + input=$COUNTERS_PATH + while IFS= read -r line || [[ -n "$line" ]]; do + #skip empty lines + if [[ -z "$line" ]]; then + continue + fi + # if in att mode, only add the first line + if [[ ! -n "$PMC_LINES" ]] || [[ ! -n "$ATT_ARGV" ]]; then + PMC_LINES+=("$line") + fi + done <$input +fi + +if [ -z $ROCPROFILER_PLUGIN_LIB ]; then + PLUGIN_TO_BE_USED=file_plugin + if [[ -z $OUTPUT_PATH && -z $OUT_FILE_NAME ]]; then + PLUGIN_TO_BE_USED=cli_plugin + fi + if [ -f $ROCM_DIR/lib/rocprofiler/$ROCPROFILER_PLUGIN_LIB$ROCPROFILER_PLUGIN_LIB_VERSION ]; then + export ROCPROFILER_PLUGIN_LIB=lib$PLUGIN_TO_BE_USED.so$ROCPROFILER_PLUGIN_LIB_VERSION + else + export ROCPROFILER_PLUGIN_LIB=lib$PLUGIN_TO_BE_USED.so + fi +fi + +export LD_PRELOAD=$LD_PRELOAD:$ROCM_DIR/lib/rocprofiler/librocprofiler_tool.so + +if [ -n "$PMC_LINES" ] && [ ! -n "$ATT_ARGV" ]; then + COUNTER=1 + + for i in ${!PMC_LINES[@]}; do + if [[ ${PMC_LINES[$i]} =~ "kernel:" ]]; then + echo ${PMC_LINES[$i]} + export ROCPROFILER_KERNEL_FILTER="${PMC_LINES[$i]}" + fi + done + + for i in ${!PMC_LINES[@]}; do + export ROCPROFILER_COUNTERS="${PMC_LINES[$i]}" + if [[ ! ${PMC_LINES[$i]} =~ "pmc" ]]; then + continue + fi + + FINAL_PATH="$OUTPUT_PATH_INTERNAL/pmc_$COUNTER" + if [ -n "$OUT_FILE_NAME" ] || [ -n "$OUTPUT_PATH" ]; then + export OUTPUT_PATH=$FINAL_PATH + fi + let COUNTER=COUNTER+1 + LD_PRELOAD=$LD_PRELOAD "${@}" + if [ -n "$OUTPUT_PATH" ]; then + echo -e "\nThe output path for the following counters: $OUTPUT_PATH" + fi + done +else + LD_PRELOAD=$LD_PRELOAD "${@}" +fi + +get_pmc_results_txt_path() { + for file_name in $(ls $1); do + if [[ $file_name == *results.txt ]]; then + echo "$1/$file_name" + fi + done +} + +if [ -n "$ATT_PATH" ]; then + if [ -n "$ATT_ARGV" ]; then + eval "$ATT_PYTHON3_ARG $ATT_PATH $ATT_ARGV" + elif [ ! -n "$PMC_LINES" ]; then + echo "ATT File is required!" + fi +fi diff --git a/projects/rocprofiler/bin/rpl_run.sh b/projects/rocprofiler/bin/rpl_run.sh new file mode 100755 index 0000000000..853fe16b30 --- /dev/null +++ b/projects/rocprofiler/bin/rpl_run.sh @@ -0,0 +1,711 @@ +#!/bin/bash + +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +ROCPROF_ARGS="$*" +time_stamp=`date +%y%m%d_%H%M%S` +BIN_DIR=$(dirname $(realpath ${BASH_SOURCE[0]})) +ROOT_DIR=$(dirname $BIN_DIR) +RUN_DIR=`pwd` +TMP_DIR="/tmp" +DATA_DIR="rpl_data_${time_stamp}_$$" + +RPL_PATH=$ROOT_DIR/lib +TLIB_PATH=$RPL_PATH/rocprofiler +TTLIB_PATH=$ROOT_DIR/lib/roctracer +ROCM_LIB_PATH=$ROOT_DIR/lib +PROF_BIN_DIR=$ROOT_DIR/libexec/rocprofiler + +# Define color code +GREEN='\033[0;32m' +GREY='\033[0;90m' +RED='\033[0;31m' +RESET='\033[0m' + +# Function to display deprecation warning message +display_warning() { + echo -e "\n${RED}WARNING: We are phasing out development and support for roctracer/rocprofiler/rocprof/rocprofv2 in favor of \ +rocprofiler-sdk/rocprofv3 in upcoming ROCm releases. Going forward, only critical defect fixes will be addressed for \ +older versions of profiling tools and libraries. We encourage all users to upgrade to the latest version, \ +rocprofiler-sdk library and rocprofv3 tool, to ensure continued support and access to new features.${RESET}\n" +} + +# Display the warning message +display_warning + +# check if rocprof is supportd on this gpu arch +V1_SUPPORTED_GPU_ARCHS=("gfx80x","gfx90x","gfx94x") +IS_SUPPORTED="false" +if [ -f $BIN_DIR/rocm_agent_enumerator ]; then + CURRENT_AGENTS_LIST=$($BIN_DIR/rocm_agent_enumerator) +else + IS_SUPPORTED="false" + echo -e "Warning: Missing rocm_agent_enumerator binary" + exit 0 +fi + +if [ -z "$ROCP_PYTHON_VERSION" ] ; then + ROCP_PYTHON_VERSION=python3 +fi + +# runtime API trace +ROCTX_TRACE=0 +HSA_TRACE=0 +SYS_TRACE=0 +HIP_TRACE=0 + +# Generate stats +GEN_STATS=0 + +# Quoting profiled cmd line +CMD_QTS=1 + +export PATH=.:$PATH + +# enable error logging +export HSA_TOOLS_REPORT_LOAD_FAILURE=1 +export HSA_VEN_AMD_AQLPROFILE_LOG=1 +export ROCPROFILER_LOG=1 +unset ROCPROFILER_SESS + +# Profiler environment +# Loading of profiler library by HSA runtime +MY_HSA_TOOLS_LIB="$RPL_PATH/librocprofiler64.so.1" + +# Tool internal Preloads +ROCPROFV1_LD_PRELOAD=$MY_HSA_TOOLS_LIB +ROCPROFV1_TOOL_PRELOAD=$MY_HSA_TOOLS_LIB + +# Loading of the test tool by ROC Profiler +export ROCP_TOOL_LIB=$TLIB_PATH/librocprof-tool.so +# ROC Profiler metrics definition +export ROCP_METRICS=$TLIB_PATH/metrics.xml + +# Enabling HSA dispatches intercepting by ROC PRofiler +export ROCP_HSA_INTERCEPT=1 +# Disabling internal ROC Profiler proxy queue (simple version supported for testing purposes) +unset ROCP_PROXY_QUEUE +# Disable AQL-profile read API +export AQLPROFILE_READ_API=0 +# ROC Profiler package path +export ROCP_PACKAGE_DIR=$ROOT_DIR +# enabled SPM KFD mode +export ROCP_SPM_KFD_MODE=1 + +# error handling +fatal() { + echo "$0: Error: $1" + echo "" + usage + exit 1 +} + +error() { + echo "$0: Error: $1" + echo "" + exit 1 +} + +error_message="" +errck() { + if [ -n "$error_message" ]; then + fatal "$1 : $error_message" + fi +} + +# usage method +usage() { + bin_name=`basename $0` + echo "ROCm Profiling Library (RPL) run script, a part of ROCprofiler library package." + echo "Full path: $BIN_DIR/$bin_name" + echo "Metrics definition: $TLIB_PATH/metrics.xml" + echo "" + echo "Usage:" + echo " $bin_name [-h] [--list-basic] [--list-derived] [-i ] [-o ] " + echo "" + echo "Options:" + echo " -h - this help" + echo " --tool-version <1|2> - to use specific version of rocprof tool, by default v1 is used" + echo " 1 - rocprofiler tool v1" + echo " 2 - rocprofiler tool v2" + echo " --verbose - verbose mode, dumping all base counters used in the input metrics" + echo " --list-basic - to print the list of basic HW counters" + echo " --list-derived - to print the list of derived metrics with formulas" + echo " --cmd-qts - quoting profiled cmd line [on]" + echo "" + echo " -i <.txt|.xml file> - input file" + echo " Input file .txt format, automatically rerun application for every profiling features line:" + echo "" + echo " # Perf counters group 1" + echo " pmc : Wavefronts VALUInsts SALUInsts SFetchInsts FlatVMemInsts LDSInsts FlatLDSInsts GDSInsts VALUUtilization FetchSize" + echo " # Perf counters group 2" + echo " pmc : WriteSize L2CacheHit" + echo " # Filter by dispatches range, GPU index and kernel names" + echo " # supported range formats: \"3:9\", \"3:\", \"3\"" + echo " range: 1 : 4" + echo " gpu: 0 1 2 3" + echo " kernel: simple Pass1 simpleConvolutionPass2" + echo "" + echo " Input file .xml format, for single profiling run:" + echo "" + echo " # Metrics list definition, also the form \":\" can be used" + echo " # All defined metrics can be found in the 'metrics.xml'" + echo " # There are basic metrics for raw HW counters and high-level metrics for derived counters" + echo " " + echo "" + echo " # Filter by dispatches range, GPU index and kernel names" + echo " " + echo "" + echo " -o - output CSV file [.csv]" + echo " -d - directory where profiler store profiling data including traces [/tmp]" + echo " The data directory is automatically removed if it is matching the default temporary directory." + echo " -t - to change the temporary directory [/tmp]" + echo " By changing the temporary directory you can prevent removing the profiling data from /tmp or enable removing from not '/tmp' directory." + echo " -m - file defining custom metrics to use in-place of defaults." + echo "" + echo " --basenames - to turn on/off truncating of the kernel full function names till the base ones [off]" + echo " --timestamp - to turn on/off the kernel dispatches timestamps, dispatch/begin/end/complete during kernel profiling [off]" + echo " --ctx-wait - to wait for outstanding contexts on profiler exit [on]" + echo " --ctx-limit - maximum number of outstanding contexts [0 - unlimited]" + echo " --heartbeat - to print progress heartbeats [0 - disabled]" + echo " --obj-tracking - to turn on/off kernels code objects tracking [on]" + echo " To support V3 code object" + echo "" + echo " --stats - generating kernel execution stats, file .stats.csv" + echo "" + echo " --roctx-trace - to enable rocTX application code annotation trace, \"Markers and Ranges\" JSON trace section." + echo " --hip-trace - to trace HIP, generates API execution stats and JSON file chrome-tracing compatible" + echo " --hsa-trace - to trace HSA, generates API execution stats and JSON file chrome-tracing compatible" + echo " --sys-trace - to trace HIP/HSA APIs and GPU activity, generates stats and JSON trace chrome-tracing compatible" + echo " '--hsa-trace' can be used in addition to select activity tracing from HSA (ROCr runtime) level" + echo " Generated files: ._stats.txt .json" + echo " Traced API list can be set by input .txt or .xml files." + echo " Input .txt:" + echo " hsa: hsa_queue_create hsa_amd_memory_pool_allocate" + echo " Input .xml:" + echo " " + echo " " + echo " " + echo " " + echo "" + echo " --roctx-rename - to rename kernels with their enclosing rocTX range's message." + echo "" + echo " --trace-start - to enable tracing on start [on]" + echo " --trace-period - to enable trace with initial delay, with periodic sample length and rate" + echo " Supported time formats: " + echo " --flush-rate - to enable trace flush rate (time period)" + echo " Supported time formats: " + echo " --parallel-kernels - to enable concurrent kernels" + echo "" + echo "Configuration file:" + echo " You can set your parameters defaults preferences in the configuration file 'rpl_rc.xml'. The search path sequence: .:${HOME}:" + echo " First the configuration file is searched in the current directory, then in the current user's home directory, and then in the installation directory." + echo " Configurable options: 'basenames', 'timestamp', 'ctx-limit', 'heartbeat', 'obj-tracking'." + echo " An example of 'rpl_rc.xml':" + echo " " + echo "" + echo " --merge-traces - Script for aggregating results from multiple rocprofiler out directries." + echo " Usage: if running with rocprof" + echo " rocprof --merge-traces -o [...]" + echo "" + exit 1 +} + +# checking for availability of rocminfo utility +if !command -v rocminfo > /dev/null 2>&1 ; then + error "'rocminfo' utility is not found: please add ROCM bin path to PATH env var."; +fi + +# setting ROCM_LIB_PATH +set_rocm_lib_path() { + + for ROCM_LIB_PATH in "$ROOT_DIR/lib" "$ROOT_DIR/lib64" ; do + if [ -d "$ROCM_LIB_PATH" ]; then + return 0 + fi + done + + #error + return 255 +} + +# profiling run method +OUTPUT_LIST="" +run() { + if ! set_rocm_lib_path ; then + echo " Fatal could not find ROCm lib directory " + fatal + fi + + # split the CURRENT_AGENTS_LIST array into individual elements. + current_gpus=(${CURRENT_AGENTS_LIST[@]}) + + for gpu in "${current_gpus[@]}"; do + + # Check first 5 characters of gpu strings. + if [[ "${V1_SUPPORTED_GPU_ARCHS[@]:0:5}" =~ "${gpu:0:5}" ]]; then + IS_SUPPORTED="true" + fi + + done + + if [[ $IS_SUPPORTED == "false" ]]; then + echo "" + echo "------------ ------------ ------------" + echo "WARNING: rocprof(v1) is not supported on this device. Recommended use: rocprofv2" + echo "Please refer project's README for a list of supported architecures." + echo "------------ ------------ ------------" + echo "" + fi + + export ROCP_INPUT="$1" + OUTPUT_DIR="$2" + shift + shift + APP_CMD=$* + + if [ "$OUTPUT_DIR" = "-" ] ; then + input_tag=`echo $ROCP_INPUT | sed "s/\.xml//"` + export ROCP_OUTPUT_DIR=${input_tag}_results_${time_stamp} + elif [ "$OUTPUT_DIR" = "--" ] ; then + unset ROCP_OUTPUT_DIR + else + export ROCP_OUTPUT_DIR=$OUTPUT_DIR + fi + echo "RPL: result dir '$ROCP_OUTPUT_DIR'" + + if [ ! -e "$ROCP_INPUT" ] ; then + error "Input file '$ROCP_INPUT' not found" + fi + + if [ -n "$ROCP_OUTPUT_DIR" ] ; then + if [ "$OUTPUT_DIR" = "-" ] ; then + if [ -e "$ROCP_OUTPUT_DIR" ] ; then + error "generated dir '$ROCP_OUTPUT_DIR' exists" + fi + fi + mkdir -p "$ROCP_OUTPUT_DIR" + OUTPUT_LIST="$OUTPUT_LIST $ROCP_OUTPUT_DIR/results.txt" + fi + + API_TRACE="" + MY_LD_PRELOAD="" + if [ "$ROCTX_TRACE" = 1 ] ; then + API_TRACE=${API_TRACE}":roctx" + MY_LD_PRELOAD="$ROCM_LIB_PATH/libroctx64.so" + fi + if [ "$HIP_TRACE" = 1 ] ; then + API_TRACE=${API_TRACE}":hip" + fi + if [ "$SYS_TRACE" = 1 ] ; then + API_TRACE=${API_TRACE}":sys" + fi + + if [ "$HSA_TRACE" = 1 ] ; then + export ROCTRACER_DOMAIN=$API_TRACE":hsa" + ROCPROFV1_TOOL_PRELOAD="$MY_LD_PRELOAD $TTLIB_PATH/libroctracer_tool.so" + ROCPROFV1_LD_PRELOAD="$MY_HSA_TOOLS_LIB:$ROCM_LIB_PATH/libroctracer64.so.4" + elif [ -n "$API_TRACE" ] ; then + export ROCTRACER_DOMAIN=$API_TRACE + OUTPUT_LIST="$ROCP_OUTPUT_DIR/" + ROCPROFV1_TOOL_PRELOAD="$MY_LD_PRELOAD $TTLIB_PATH/libroctracer_tool.so" + ROCPROFV1_LD_PRELOAD="$ROCM_LIB_PATH/libroctracer64.so.4" + fi + + if [ "$ROCP_STATS_OPT" = 1 ] ; then + if [ "$ROCTRACER_DOMAIN" = ":hip" ] ; then + ROCPROFV1_LD_PRELOAD="$ROCM_LIB_PATH/libroctracer64.so.4" + ROCPROFV1_TOOL_PRELOAD="$MY_LD_PRELOAD $TTLIB_PATH/libhip_stats.so" + else + error_message="ROCP_STATS_OPT is only available with --hip-trace option" + echo $error_message + exit 1 + fi + fi + + retval=1 + if [ -n "$ROCP_OUTPUT_DIR" ] ; then + log_file="$ROCP_OUTPUT_DIR/log.txt" + exit_file="$ROCP_OUTPUT_DIR/exit.txt" + { + HSA_TOOLS_LIB="$ROCPROFV1_TOOL_PRELOAD" LD_PRELOAD=$LD_PRELOAD:"$ROCPROFV1_LD_PRELOAD" eval "$APP_CMD" + retval=$? + echo "exit($retval)" > $exit_file + } 2>&1 | tee "$log_file" + exitval=`cat "$exit_file" | sed -n "s/^.*exit(\([0-9]*\)).*$/\1/p"` + if [ -n "$exitval" ] ; then retval=$exitval; fi + else + HSA_TOOLS_LIB="$ROCPROFV1_TOOL_PRELOAD" LD_PRELOAD=$LD_PRELOAD:"$ROCPROFV1_LD_PRELOAD" eval "$APP_CMD" + retval=$? + fi + return $retval +} + +merge_output() { + while [ -n "$1" ] ; do + output_dir=$(echo "$1" | sed "s/\/[^\/]*$//") + for file_name in `ls $output_dir` ; do + output_name=$(echo $file_name | sed -n "/\.txt$/ s/^[0-9]*_//p") + if [ -n "$output_name" ] ; then + trace_file=$output_dir/$file_name + output_file=$output_dir/$output_name + touch $output_file + cat $trace_file >> $output_file + fi + done + shift + done +} + +convert_time_val() { + local time_maxumim_us=$((0xffffffff)) + local __resultvar=$1 + eval "local val=$"$__resultvar + val_m=`echo $val | sed -n "s/^\([0-9]*\)m$/\1/p"` + val_s=`echo $val | sed -n "s/^\([0-9]*\)s$/\1/p"` + val_ms=`echo $val | sed -n "s/^\([0-9]*\)ms$/\1/p"` + val_us=`echo $val | sed -n "s/^\([0-9]*\)us$/\1/p"` + if [ -n "$val_m" ] ; then val_us=$((val_m*60000000)) + elif [ -n "$val_s" ] ; then val_us=$((val_s*1000000)) + elif [ -n "$val_ms" ] ; then val_us=$((val_ms*1000)) + fi + + if [ -z "$val_us" ] ; then + error_message="invalid time value format ($val)" + elif [ "$val_us" -gt "$time_maxumim_us" ] ; then + error_message="time value exceeds maximum supported ($val > ${time_maxumim_us}us)" + else + eval $__resultvar="'$val_us'" + fi +} + +unset_v1_envs() { + # enable error logging + unset HSA_TOOLS_REPORT_LOAD_FAILURE + unset HSA_VEN_AMD_AQLPROFILE_LOG + unset ROCPROFILER_LOG + unset ROCPROFILER_SESS + unset MY_HSA_TOOLS_LIB + unset ROCP_TOOL_LIB + unset ROCP_METRICS + unset ROCP_HSA_INTERCEPT + unset ROCP_PROXY_QUEUE + unset AQLPROFILE_READ_API + unset ROCP_PACKAGE_DIR + unset ROCP_SPM_KFD_MODE +} + +################################################################################################ +# main +echo "RPL: on '$time_stamp' from '$ROOT_DIR' in '$RUN_DIR'" +# Parsing arguments +if [ -z "$1" ] ; then + usage +fi + +INPUT_FILE="" +DATA_PATH="-" +OUTPUT_DIR="-" +output="" +csv_output="" + +ARG_IN="" +while [ 1 ] ; do + ARG_IN=$1 + ARG_VAL=1 + if [ "$1" = "--tool-version" ] ; then + if [ $2 = 1 ] ; then + : + elif [ $2 = 2 ] ; then + # unset all v1 variables + unset_v1_envs + eval $BIN_DIR/rocprofv2 $ROCPROF_ARGS + exit 0 + else + echo "Wrong option '$1 $2'" + usage + fi + elif [ "$1" = "--version" ]; then + if [ -f "$BIN_DIR/../libexec/rocprofiler/rocprofiler-version" ]; then + ROCPROFILER_LIBRARY_VERSION=1 $BIN_DIR/../libexec/rocprofiler/rocprofiler-version + else + ROCM_VERSION=$(cat $BIN_DIR/../.info/version) + echo -e "ROCm version: $ROCM_VERSION" + echo -e "ROCProfiler version: 2.0" + fi + exit 0 + elif [ "$1" = "-h" ] ; then + usage + elif [ "$1" = "-i" ] ; then + INPUT_FILE="$2" + elif [ "$1" = "-o" ] ; then + output="$2" + if [[ $output != *.csv ]]; then + echo "error: file name must have .CSV extension" + exit 1 + fi + elif [ "$1" = "-d" ] ; then + DATA_PATH=$2 + elif [ "$1" = "-t" ] ; then + TMP_DIR="$2" + if [ "$OUTPUT_DIR" = "-" ] ; then + DATA_PATH=$TMP_DIR + fi + elif [ "$1" = "-m" ] ; then + unset ROCP_METRICS + export ROCP_METRICS="$2" + elif [ "$1" = "--list-basic" ] ; then + export ROCP_INFO=b + HSA_TOOLS_LIB="$MY_HSA_TOOLS_LIB" eval "$TLIB_PATH/rocprof-ctrl" + exit 1 + elif [ "$1" = "--list-derived" ] ; then + export ROCP_INFO=d + HSA_TOOLS_LIB="$MY_HSA_TOOLS_LIB" eval "$TLIB_PATH/rocprof-ctrl" + exit 1 + elif [ "$1" = "--basenames" ] ; then + if [ "$2" = "on" ] ; then + export ROCP_TRUNCATE_NAMES=1 + else + export ROCP_TRUNCATE_NAMES=0 + fi + elif [ "$1" = "--timestamp" ] ; then + if [ "$2" = "on" ] ; then + export ROCP_TIMESTAMP_ON=1 + else + export ROCP_TIMESTAMP_ON=0 + fi + elif [ "$1" = "--ctx-wait" ] ; then + if [ "$2" = "on" ] ; then + export ROCP_OUTSTANDING_WAIT=1 + else + export ROCP_OUTSTANDING_WAIT=0 + fi + elif [ "$1" = "--ctx-limit" ] ; then + export ROCP_OUTSTANDING_MAX="$2" + elif [ "$1" = "--heartbeat" ] ; then + export ROCP_OUTSTANDING_MON="$2" + elif [ "$1" = "--stats" ] ; then + ARG_VAL=0 + export ROCP_TIMESTAMP_ON=1 + GEN_STATS=1 + elif [ "$1" = "--roctx-trace" ] ; then + ARG_VAL=0 + GEN_STATS=1 + ROCTX_TRACE=1 + elif [ "$1" = "--hsa-trace" ] ; then + ARG_VAL=0 + export ROCP_TIMESTAMP_ON=1 + GEN_STATS=1 + HSA_TRACE=1 + elif [ "$1" = "--sys-trace" ] ; then + ARG_VAL=0 + export ROCP_TIMESTAMP_ON=1 + GEN_STATS=1 + SYS_TRACE=1 + elif [ "$1" = "--hip-trace" ] ; then + ARG_VAL=0 + export ROCP_TIMESTAMP_ON=1 + GEN_STATS=1 + HIP_TRACE=1 + elif [ "$1" = "--roctx-rename" ] ; then + ARG_VAL=0 + export ROCP_RENAME_KERNEL=1 + elif [ "$1" = "--trace-start" ] ; then + if [ "$2" = "off" ] ; then + export ROCP_CTRL_RATE="-1:0:0" + fi + elif [ "$1" = "--trace-period" ] ; then + period_expr="^\([^:]*\):\([^:]*\):\([^:]*\)$" + period_ck=`echo "$2" | sed -n "s/"${period_expr}"/ok/p"` + if [ -z "$period_ck" ] ; then + fatal "Wrong option '$1 $2'" + fi + period_delay=`echo "$2" | sed -n "s/"${period_expr}"/\1/p"` + period_len=`echo "$2" | sed -n "s/"${period_expr}"/\2/p"` + period_rate=`echo "$2" | sed -n "s/"${period_expr}"/\3/p"` + convert_time_val period_delay + errck "Option '$ARG_IN', delay value" + convert_time_val period_len + errck "Option '$ARG_IN', length value" + convert_time_val period_rate + errck "Option '$ARG_IN', rate value" + export ROCP_CTRL_RATE="$period_delay:$period_len:$period_rate" + elif [ "$1" = "--flush-rate" ] ; then + period_rate=$2 + convert_time_val period_rate + errck "Option '$ARG_IN', rate value" + export ROCP_FLUSH_RATE="$period_rate" + elif [ "$1" = "--obj-tracking" ] ; then + if [ "$2" = "off" ] ; then + export ROCP_OBJ_TRACKING=0 + fi + elif [ "$1" = "--parallel-kernels" ] ; then + ARG_VAL=0 + export ROCP_K_CONCURRENT=1 + export AQLPROFILE_READ_API=1 + elif [ "$1" = "--verbose" ] ; then + ARG_VAL=0 + export ROCP_VERBOSE_MODE=1 + elif [ "$1" = "--cmd-qts" ] ; then + if [ "$2" = "off" ] ; then + CMD_QTS=0 + fi + elif [ "$1" = "--merge-traces" ] ; then + shift + echo "merging traces with $PROF_BIN_DIR/merge_traces.sh" + $PROF_BIN_DIR/merge_traces.sh $@ + exit 0 + else + break + fi + shift + if [ "$ARG_VAL" = 1 ] ; then shift; fi +done + +ARG_CK=`echo $ARG_IN | sed "s/^-.*$/-/"` +if [ "$ARG_CK" = "-" ] ; then + fatal "Wrong option '$ARG_IN'" +fi + +if [ "$GEN_STATS" = "1" -a "$ROCP_TIMESTAMP_ON" = "0" ] ; then + fatal "Wrong options, stats enabled with disabled timestamps" +fi + +if [ -z "$INPUT_FILE" ] ; then + input_base="results" + input_type="none" +else + input_base=`echo "$INPUT_FILE" | sed "s/^\(.*\)\.\([^\.]*\)$/\1/"` + input_type=`echo "$INPUT_FILE" | sed "s/^\(.*\)\.\([^\.]*\)$/\2/"` + if [ -z "${input_base}" -o -z "${input_type}" ] ; then + fatal "Bad input file '$INPUT_FILE'" + fi + input_base=`basename $input_base` +fi + +if [ "$DATA_PATH" = "-" ] ; then + DATA_PATH=$TMP_DIR +fi + +if [ -n "$output" ] ; then + if [ "$output" = "--" ] ; then + OUTPUT_DIR="--" + else + csv_output=$output + fi +else + csv_output=$RUN_DIR/${input_base}.csv +fi + +# Profiled cmd line string +APP_CMD=$* +if [ "$CMD_QTS" = 1 ] ; then + APP_CMD="" + for i in `seq 1 $#`; do + if [ -n "$APP_CMD" ] ; then + APP_CMD=$APP_CMD" " + fi + eval "arg=\${$i}" + APP_CMD=$APP_CMD\"$arg\" + done +fi + +echo "RPL: profiling '$APP_CMD'" +echo "RPL: input file '$INPUT_FILE'" + +input_list="" +RES_DIR="" +if [ "$input_type" = "xml" ] ; then + OUTPUT_DIR=$DATA_PATH + input_list=$INPUT_FILE +elif [ "$input_type" = "txt" -o "$input_type" = "none" ] ; then + RES_DIR=$DATA_PATH/$DATA_DIR + if [ -e $RES_DIR ] ; then + error "Rundir '$RES_DIR' exists" + fi + mkdir -p $RES_DIR + echo "RPL: output dir '$RES_DIR'" + if [ "$input_type" = "txt" ] ; then + $PROF_BIN_DIR/txt2xml.sh $INPUT_FILE $RES_DIR + else + echo "" > $RES_DIR/input.xml + fi + input_list=`/bin/ls $RES_DIR/input*.xml` + export ROCPROFILER_SESS=$RES_DIR +else + fatal "Bad input file type '$INPUT_FILE'" +fi + +if [ -n "$csv_output" ] ; then + rm -f $csv_output +fi + +RET=1 +export ROCP_MERGE_PIDS=1 +for name in $input_list; do + run $name $OUTPUT_DIR $APP_CMD + RET=$? + if [ -n "$ROCPROFILER_SESS" -a -e "$ROCPROFILER_SESS/error" ] ; then + error_string=`cat $ROCPROFILER_SESS/error` + echo "Profiling error found: '$error_string'" + csv_output="" + RET=1 + break + fi +done + +if [ -n "$csv_output" ] ; then + merge_output $OUTPUT_LIST + if [ "$GEN_STATS" = "1" ] ; then + db_output=$(echo $csv_output | sed "s/\.csv/.db/") + $ROCP_PYTHON_VERSION $PROF_BIN_DIR/tblextr.py $db_output $OUTPUT_LIST + else + $ROCP_PYTHON_VERSION $PROF_BIN_DIR/tblextr.py $csv_output $OUTPUT_LIST + fi + if [ "$?" -ne 0 ] ; then + echo "Profiling data corrupted: '$OUTPUT_LIST'" | tee "$ROCPROFILER_SESS/error" + RET=1 + fi +fi + +if [ "$DATA_PATH" = "$TMP_DIR" ] ; then + if [ -e "$RES_DIR" ] ; then + rm -rf $RES_DIR + fi +fi + +exit $RET diff --git a/projects/rocprofiler/bin/run_tool.sh b/projects/rocprofiler/bin/run_tool.sh new file mode 100755 index 0000000000..b61db70433 --- /dev/null +++ b/projects/rocprofiler/bin/run_tool.sh @@ -0,0 +1,34 @@ +#!/bin/sh +BIN_DIR=`dirname $0` +BIN_DIR=`realpath $BIN_DIR` +PKG_DIR=${BIN_DIR%/bin} + +if [ -z "$1" ] ; then + echo "Usage: $0 " + exit 1 +fi + +# profiler plugin library +test_app=$* + +# ROC profiler library loaded by HSA runtime +export HSA_TOOLS_LIB=librocprofiler64.so.1 + +# tool library loaded by ROC profiler +if [ -z "$ROCP_TOOL_LIB" ] ; then + echo "ROCP_TOOL_LIB is not found" + exit 1 +fi + +# ROC profiler metrics config file +if [ -z "$ROCP_METRICS" ] ; then + echo "ROCP_METRICS is not found" + exit 1 +fi + +# enable error messages +export HSA_TOOLS_REPORT_LOAD_FAILURE=1 +export HSA_VEN_AMD_AQLPROFILE_LOG=1 +export ROCPROFILER_LOG=1 + +$test_app diff --git a/projects/rocprofiler/bin/sqlitedb.py b/projects/rocprofiler/bin/sqlitedb.py new file mode 100644 index 0000000000..4d54e29a0b --- /dev/null +++ b/projects/rocprofiler/bin/sqlitedb.py @@ -0,0 +1,351 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +import csv, sqlite3, re, sys +from functools import reduce +from txt2params import gen_params + + +# SQLite Database class +class SQLiteDB: + def __init__(self, file_name): + self.connection = sqlite3.connect(file_name) + self.tables = {} + self.section_index = 0 + + def __del__(self): + self.connection.close() + + # add DB table + def add_table(self, name, descr, extra=()): + (field_list, field_dict) = descr + if name in self.tables: + raise Exception('table is already added: "' + name + '"') + + # create DB table + table_descr = [] + for field in field_list: + table_descr.append('"%s" %s' % (field, field_dict[field])) + for item in extra: + table_descr.append('"%s" %s' % (item[0], item[1])) + stm = "CREATE TABLE " + name + " (%s)" % ", ".join(table_descr) + cursor = self.connection.cursor() + cursor.execute(stm) + self.connection.commit() + + # register table + fields_str = ",".join(map(lambda x: '"' + x + '"', field_list)) + templ_str = ",".join("?" * len(field_list)) + stm = "INSERT INTO " + name + "(" + fields_str + ") VALUES(" + templ_str + ");" + self.tables[name] = stm + + return (cursor, stm) + + # add columns to table + def add_columns(self, name, columns): + cursor = self.connection.cursor() + for item in columns: + stm = "ALTER TABLE " + name + ' ADD COLUMN "%s" %s' % (item[0], item[1]) + cursor.execute(stm) + self.connection.commit() + + # add columns with expression + def add_data_column(self, table_name, data_label, data_type, data_expr): + cursor = self.connection.cursor() + cursor.execute( + 'ALTER TABLE %s ADD COLUMN "%s" %s' % (table_name, data_label, data_type) + ) + cursor.execute("UPDATE %s SET %s = (%s);" % (table_name, data_label, data_expr)) + + def change_rec_name(self, table_name, rec_id, rec_name): + self.connection.execute( + "UPDATE " + table_name + ' SET Name = ? WHERE "Index" = ?', (rec_name, rec_id) + ) + + def change_rec_tid(self, table_name, rec_id, tid): + self.connection.execute( + "UPDATE " + table_name + ' SET tid = ? WHERE "Index" = ?', (tid, rec_id) + ) + + def change_rec_fld(self, table_name, fld_expr, rec_pat): + self.connection.execute( + "UPDATE " + table_name + " SET " + fld_expr + " WHERE " + rec_pat + ) + + def table_get_record(self, table_name, rec_pat): + cursor = self.connection.execute( + "SELECT * FROM " + table_name + " WHERE " + rec_pat + ) + raws = cursor.fetchall() + if len(raws) != 1: + raise Exception( + "Record (" + rec_pat + ') is not unique, table "' + table_name + '"' + ) + return list(raws[0]) + + # populate DB table entry + def insert_entry(self, table, val_list): + (cursor, stm) = table + cursor.execute(stm, val_list) + + # populate DB table entry + def commit_entry(self, table, val_list): + self.insert_entry(table, val_list) + self.connection.commit() + + # populate DB table data + def insert_table(self, table, reader): + for val_list in reader: + if not val_list[-1]: + val_list.pop() + self.insert_entry(table, val_list) + self.connection.commit() + + # return table fields list + def _get_fields(self, table_name): + cursor = self.connection.execute("SELECT * FROM " + table_name) + return list(map(lambda x: '"%s"' % (x[0]), cursor.description)) + + # return table raws list + def _get_raws(self, table_name): + cursor = self.connection.execute("SELECT * FROM " + table_name) + return cursor.fetchall() + + def _get_raws_indexed(self, table_name): + cursor = self.connection.execute( + "SELECT * FROM " + table_name + ' order by "Index" asc;' + ) + return cursor.fetchall() + + def _get_raw_by_id(self, table_name, rec_id): + cursor = self.connection.execute( + "SELECT * FROM " + table_name + ' WHERE "Index"=?', (rec_id,) + ) + raws = cursor.fetchall() + if len(raws) != 1: + raise Exception('Index is not unique, table "' + table_name + '"') + return list(raws[0]) + + def table_get_raws(self, table_name): + return self._get_raws(table_name) + + # dump CSV table + def dump_csv(self, table_name, file_name): + if not re.search(r"\.csv$", file_name): + raise Exception('wrong output file type: "' + file_name + '"') + + fields = self._get_fields(table_name) + with open(file_name, mode="w") as fd: + fd.write(",".join(fields) + "\n") + for raw in self._get_raws(table_name): + tmp = list(raw) + for idx in range(len(tmp)): + if type(tmp[idx]) == str: + if not (tmp[idx][0] == tmp[idx][-1] == '"'): + tmp[idx] = '"' + tmp[idx] + '"' + raw = tuple(tmp) + fd.write(reduce(lambda a, b: str(a) + "," + str(b), raw) + "\n") + + # dump JSON trace + def open_json(self, file_name): + if not re.search(r"\.json$", file_name): + raise Exception('wrong output file type: "' + file_name + '"') + with open(file_name, mode="w") as fd: + fd.write('{ "traceEvents":[{}\n') + + def close_json(self, file_name): + if not re.search(r"\.json$", file_name): + raise Exception('wrong output file type: "' + file_name + '"') + with open(file_name, mode="a") as fd: + fd.write("}") + + def label_json(self, pid, label, file_name): + if not re.search(r"\.json$", file_name): + raise Exception('wrong output file type: "' + file_name + '"') + with open(file_name, mode="a") as fd: + fd.write( + ',{"args":{"name":"%s"},"ph":"M","pid":%s,"name":"process_name","sort_index":%d}\n' + % (label, pid, self.section_index) + ) + self.section_index += 1 + + def flow_json( + self, base_id, from_pid, from_us_list, to_pid, to_us_dict, corr_id_list, file_name + ): + if not re.search(r"\.json$", file_name): + raise Exception('wrong output file type: "' + file_name + '"') + with open(file_name, mode="a") as fd: + dep_id = base_id + for ind in range(len(from_us_list)): + corr_id = corr_id_list[ind] if (len(corr_id_list) != 0) else ind + if corr_id in to_us_dict: + (from_ts, stream_id, tid) = from_us_list[ind] + to_ts = to_us_dict[corr_id] + if from_ts > to_ts: + from_ts = to_ts + fd.write( + ',{"ts":%d,"ph":"s","cat":"DataFlow","id":%d,"pid":%d,"tid":%d,"name":"dep"}\n' + % (from_ts, dep_id, from_pid, tid) + ) + fd.write( + ',{"ts":%d,"ph":"t","cat":"DataFlow","id":%d,"pid":%d,"tid":%d,"name":"dep"}\n' + % (to_ts, dep_id, to_pid, stream_id) + ) + dep_id += 1 + + def metadata_json(self, jsonfile, sysinfo_file): + params = gen_params(sysinfo_file) + with open(jsonfile, mode="a") as fd: + cnt = 0 + fd.write("],\n") + fd.write('"otherData": {\n') + for nkey in sorted(params.keys()): + key = nkey[1] + cnt = cnt + 1 + if cnt == len(params): + fd.write(' "' + key + '": "' + params[nkey] + '"\n') + else: + fd.write(' "' + key + '": "' + params[nkey] + '",\n') + fd.write(" }\n") + + def dump_json(self, table_name, data_name, file_name): + if not re.search(r"\.json$", file_name): + raise Exception('wrong output file type: "' + file_name + '"') + + sub_ptrn = re.compile(r'(^"|"$)') + name_ptrn = re.compile(r"(name|Name)") + + table_fields = self._get_fields(table_name) + table_raws = self._get_raws(table_name) + data_fields = self._get_fields(data_name) + data_raws = self._get_raws(data_name) + + with open(file_name, mode="a") as fd: + table_raws_len = len(table_raws) + for raw_index in range(table_raws_len): + if (raw_index == table_raws_len - 1) or (raw_index % 1000 == 0): + sys.stdout.write( + "\rdump json " + + str(raw_index) + + ":" + + str(len(table_raws)) + + " " * 100 + ) + + vals_list = [] + values = list(table_raws[raw_index]) + for value_index in range(len(values)): + label = table_fields[value_index] + value = values[value_index] + if name_ptrn.search(label): + value = sub_ptrn.sub(r"", value) + if label != '"Index"': + if label == '"dur"' and value == 0: + vals_list.append('%s:"%s"' % (label, "1")) + elif label == '"pid"' or label == '"tid"': + vals_list.append('%s:%d' % (label, value)) + else: + vals_list.append('%s:"%s"' % (label, value)) + + args_list = [] + data = list(data_raws[raw_index]) + for value_index in range(len(data)): + label = data_fields[value_index] + value = data[value_index] + if label[:3] == '"__': + continue + if name_ptrn.search(label): + value = sub_ptrn.sub(r"", value) + if label != '"Index"' and label != '"roctx-range"': + args_list.append('%s:"%s"' % (label, value)) + + fd.write( + ',{"ph":"%s",%s,\n "args":{\n %s\n }\n}\n' + % ("X", ",".join(vals_list), ",\n ".join(args_list)) + ) + + sys.stdout.write("\n") + + # execute query on DB + def execute(self, cmd): + cursor = self.connection.cursor() + cursor.execute(cmd) + + # commit DB + def commit(self): + self.connection.commit() + + # close DB + def close(self): + self.connection.commit() + self.connection.close() + + # access DB + def get_raws(self, table_name): + cur = self.connection.cursor() + cur.execute("SELECT * FROM %s" % table_name) + return cur.fetchall() + + # return CSV descriptor + # list of fields and dictionaly for the fields types + def _get_csv_descr(self, table_name, fd): + reader = csv.DictReader(fd) + field_names = reader.fieldnames + if not field_names[-1]: + field_names.pop() + field_types = {} + + for entry in reader: + fields_left = [f for f in field_names if f not in field_types.keys()] + # all fields processed + if not fields_left: + break + + for field in fields_left: + data = entry[field] + # need data for the field to be processed + if len(data) == 0: + continue + + if data.isdigit(): + field_types[field] = "INTEGER" + else: + field_types[field] = "TEXT" + + if len(fields_left) > 0: + raise Exception("types not found for fields: ", fields_left) + return (field_names, field_types) + + # add CSV table + def add_csv_table(self, table_name, file_name, extra=()): + with open(file_name, mode="r") as fd: + # get CSV table descriptor + descr = self._get_csv_descr(table_name, fd) + # reader to populate the table + fd.seek(0) + reader = csv.reader(fd) + reader.next() + table = self.add_table(table_name, descr, extra) + self.insert_table(table, reader) + + +############################################################################################## diff --git a/projects/rocprofiler/bin/tblextr.py b/projects/rocprofiler/bin/tblextr.py new file mode 100755 index 0000000000..f9295e05eb --- /dev/null +++ b/projects/rocprofiler/bin/tblextr.py @@ -0,0 +1,1242 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +import os +import sys +import re +import subprocess +import bisect +from sqlitedb import SQLiteDB +from mem_manager import MemManager +import dform + +mcopy_data_enabled = 0 + +EXT_PID = 0 +COPY_PID = 1 +HIP_PID = 2 +HSA_PID = 3 +OPS_PID = 5 +GPU_BASE_PID = 6 +NONE_PID = -1 + +max_gpu_id = -1 +START_NS = 0 + +hsa_activity_found = 0 + +# dependencies dictionary +dep_dict = {} +kern_dep_list = [] +last_hip_api_map = {} +hip_streams = [] +from_ids = {} + +# stream ID map +stream_counter = 0 +stream_id_map = {} + + +def get_stream_index(stream_id): + global stream_counter + stream_ind = 0 + if stream_id.lower() != "nil": + if not stream_id in stream_id_map: + stream_counter += 1 + stream_ind = stream_counter + stream_id_map[stream_id] = stream_ind + else: + stream_ind = stream_id_map[stream_id] + return stream_ind + + +# patching activity records +def activity_record_patching( + db, ops_table_name, kernel_found, kernel_name, stream_found, stream_ind, select_expr +): + if kernel_found != 0: + db.change_rec_fld(ops_table_name, 'Name = "' + kernel_name + '"', select_expr) + if stream_found != 0: + db.change_rec_fld(ops_table_name, "tid = " + str(stream_ind), select_expr) + + +# global vars +table_descr = [["Index", "KernelName"], {"Index": "INTEGER", "KernelName": "TEXT"}] +var_list = table_descr[0] +var_table = {} +############################################################# + + +def fatal(msg): + sys.stderr.write(sys.argv[0] + ": " + msg + "\n") + sys.exit(1) + + +dbglog_count = 0 + + +def dbglog(msg): + global dbglog_count + dbglog_count += 1 + sys.stderr.write(sys.argv[0] + ": " + msg + "\n") + fatal("error") + + +############################################################# + +# Dumping sysinfo +sysinfo_begin = 1 + + +def metadata_gen(sysinfo_file, sysinfo_cmd): + global sysinfo_begin + if not re.search(r"\.txt$", sysinfo_file): + raise Exception('wrong output file type: "' + sysinfo_file + '"') + if sysinfo_begin == 1: + sysinfo_begin = 0 + with open(sysinfo_file, mode="w") as fd: + fd.write("") + with open(sysinfo_file, mode="a") as fd: + fd.write("CMD: " + sysinfo_cmd + "\n") + status = subprocess.call( + sysinfo_cmd + " >> " + sysinfo_file, stderr=subprocess.STDOUT, shell=True + ) + if status != 0: + raise Exception('Could not run command: "' + sysinfo_cmd + '"') + + +# parse results method +def parse_res(infile): + global max_gpu_id + if not os.path.isfile(infile): + return + inp = open(infile, "r", errors="replace") + + beg_pattern = re.compile(r'^dispatch\[(\d*)\], (.*) kernel-name\("([^"]*)"\)') + prop_pattern = re.compile(r"([\w-]+)\((\w+)\)") + ts_pattern = re.compile(r", time\((\d*),(\d*),(\d*),(\d*)\)") + # var pattern below matches a variable name and a variable value from a one + # line text in the format of for example "WRITE_SIZE (0.2500000000)" or + # "GRBM_GUI_ACTIVE (27867)" or "TA_TA_BUSY[0]" + var_pattern = re.compile(r"^\s*([a-zA-Z0-9_]+(?:\[\d+\])?)\s+\((\d+(?:\.\d+)?)\)") + pid_pattern = re.compile(r"pid\((\d*)\)") + + dispatch_number = 0 + var_table_pid = 0 + for line in inp.readlines(): + record = line[:-1] + + m = pid_pattern.search(record) + if m and not os.getenv("ROCP_MERGE_PIDS"): + var_table_pid = int(m.group(1)) + + m = var_pattern.match(record) + if m: + if not (var_table_pid, dispatch_number) in var_table: + fatal("Error: dispatch number not found '" + str(dispatch_number) + "'") + var = m.group(1) + val = m.group(2) + var_table[(var_table_pid, dispatch_number)][var] = val + if not var in var_list: + var_list.append(var) + + m = beg_pattern.match(record) + if m: + dispatch_number = m.group(1) + if not (var_table_pid, dispatch_number) in var_table: + var_table[(var_table_pid, dispatch_number)] = { + "Index": dispatch_number, + "KernelName": '"' + m.group(3) + '"', + } + + gpu_id = 0 + queue_id = 0 + disp_pid = 0 + disp_tid = 0 + + kernel_properties = m.group(2) + for prop in kernel_properties.split(", "): + m = prop_pattern.match(prop) + if m: + var = m.group(1) + val = m.group(2) + var_table[(var_table_pid, dispatch_number)][var] = val + if not var in var_list: + var_list.append(var) + if var == "gpu-id": + gpu_id = int(val) + if gpu_id > max_gpu_id: + max_gpu_id = gpu_id + if var == "queue-id": + queue_id = int(val) + if var == "pid": + disp_pid = int(val) + if var == "tid": + disp_tid = int(val) + else: + fatal( + 'wrong kernel property "' + + prop + + '" in "' + + kernel_properties + + '"' + ) + m = ts_pattern.search(record) + if m: + var_table[(var_table_pid, dispatch_number)]["DispatchNs"] = m.group(1) + var_table[(var_table_pid, dispatch_number)]["BeginNs"] = m.group(2) + var_table[(var_table_pid, dispatch_number)]["EndNs"] = m.group(3) + var_table[(var_table_pid, dispatch_number)]["CompleteNs"] = m.group(4) + + # filling dependenciws + from_ns = int(m.group(1)) + to_ns = int(m.group(2)) + from_us = int((from_ns - START_NS) / 1000) + to_us = int((to_ns - START_NS) / 1000) + + kern_dep_list.append((from_ns, disp_pid, disp_tid)) + + gpu_pid = GPU_BASE_PID + int(gpu_id) + if not disp_pid in dep_dict: + dep_dict[disp_pid] = {} + dep_proc = dep_dict[disp_pid] + if not gpu_pid in dep_proc: + dep_proc[gpu_pid] = { + "pid": HSA_PID, + "from": [], + "to": {}, + "id": [], + } + dep_str = dep_proc[gpu_pid] + to_id = len(dep_str["from"]) + dep_str["from"].append((from_us, disp_tid, disp_tid)) + dep_str["to"][to_id] = to_us + ## + + inp.close() + + +############################################################# + + +# Comparator to sort a dictionary of tuples. This comparator will convert +# the second element of tuple to an int and return the new tuple. Then +# the dictionary can use the default comparison i.e sort by first element, +# then sort by second element. +def tuple_comparator(tupleElem): + return tupleElem[0], int(tupleElem[1]) + + +# merge results table +def merge_table(): + global var_list + keys = sorted(var_table.keys(), key=tuple_comparator) + + fields = set(var_table[keys[0]]) + if "DispatchNs" in fields: + var_list.append("DispatchNs") + var_list.append("BeginNs") + var_list.append("EndNs") + var_list.append("CompleteNs") + var_list = [x for x in var_list if x in fields] + + +############################################################# + + +# dump CSV results +def dump_csv(file_name): + global var_list + keys = sorted(var_table.keys(), key=tuple_comparator) + + with open(file_name, mode="w") as fd: + fd.write(",".join(var_list) + "\n") + for pid, ind in keys: + entry = var_table[(pid, ind)] + dispatch_number = entry["Index"] + if ind != dispatch_number: + fatal("Dispatch #" + ind + " index mismatch (" + dispatch_number + ")\n") + val_list = [entry[var] if (var in entry) else 'None' for var in var_list] + fd.write(",".join(val_list) + "\n") + + print("File '" + file_name + "' is generating") + + +############################################################# + + +# fill kernels DB +def fill_kernel_db(table_name, db): + global var_list + keys = sorted(var_table.keys(), key=tuple_comparator) + + for var in set(var_list).difference(set(table_descr[1])): + table_descr[1][var] = "INTEGER" + table_descr[0] = var_list + + table_handle = db.add_table(table_name, table_descr) + + for pid, ind in keys: + entry = var_table[(pid, ind)] + dispatch_number = entry["Index"] + if ind != dispatch_number: + fatal("Dispatch #" + ind + " index mismatch (" + dispatch_number + ")\n") + val_list = [entry[var] for var in var_list] + db.insert_entry(table_handle, val_list) + + +############################################################# + +# Fill Ext DB +ext_table_descr = [ + ["BeginNs", "EndNs", "pid", "tid", "Name", "Index", "__section", "__lane"], + { + "BeginNs": "INTEGER", + "EndNs": "INTEGER", + "pid": "INTEGER", + "tid": "INTEGER", + "Name": "TEXT", + "Index": "INTEGER", + "__section": "INTEGER", + "__lane": "INTEGER", + }, +] + + +def fill_ext_db(table_name, db, indir, trace_name, api_pid): + global range_data + + file_name = indir + "/" + trace_name + "_trace.txt" + # tms pid:tid cid:rid:'.....' + ptrn_val = re.compile(r'(\d+) (\d+):(\d+) (\d+):(\d+):"(.*)"$') + + range_data = {} + range_stack = {} + range_map = {} + + if not os.path.isfile(file_name): + return 0 + + record_id = 0 + table_handle = db.add_table(table_name, ext_table_descr) + with open(file_name, mode="r", errors="replace") as fd: + for line in fd.readlines(): + record = line[:-1] + m = ptrn_val.match(record) + if m: + tms = int(m.group(1)) + pid = m.group(2) + tid = int(m.group(3)) + cid = int(m.group(4)) + rid = int(m.group(5)) + msg = m.group(6) + + rec_vals = [] + if not tid in range_data: + range_data[tid] = {} + + if cid != 2: + rec_vals.append(tms) + rec_vals.append(tms + 1) + rec_vals.append(pid) + rec_vals.append(tid) + rec_vals.append(msg) + rec_vals.append(record_id) + rec_vals.append(api_pid) # __section + rec_vals.append(tid) # __lane + + if cid == 1: + if not pid in range_stack: + range_stack[pid] = {} + pid_stack = range_stack[pid] + if not tid in pid_stack: + pid_stack[tid] = [] + rec_stack = pid_stack[tid] + rec_stack.append(rec_vals) + continue + + if cid == 2: + if not pid in range_stack: + fatal("ROCTX range begin not found, pid(" + pid + ")") + pid_stack = range_stack[pid] + if not tid in pid_stack: + fatal("ROCTX range begin not found, tid(" + tid + ")") + rec_stack = pid_stack[tid] + rec_vals = rec_stack.pop() + rec_vals[1] = tms + # record the range's start/stop timestamps, its parent (ranges can be nested), and its message. + range_start = rec_vals[0] + range_stop = tms + range_parent = rec_stack[-1][0] if len(rec_stack) != 0 else 0 + range_msg = rec_vals[4] + range_data[tid][range_start] = (range_stop, range_parent, range_msg) + + # range start + if cid == 3: + range_map[rid] = (tms, msg) + continue + + # range stop + if cid == 4: + if rid in range_map: + # querying start timestamp if rid exists + (tms, msg) = range_map[rid] + del range_map[rid] + else: + fatal("range id(" + str(rid) + ") is not found") + rec_vals[0] = tms # begin timestamp + rec_vals[4] = msg # range message + rec_vals[7] = 0 # 0 lane for ranges + + db.insert_entry(table_handle, rec_vals) + record_id += 1 + + return 1 + + +############################################################# +# arguments manipulation routines +def get_field(args, field): + if args == None: + return (None, 0) + ptrn1_field = re.compile(r"^.* " + field + r"\(") + ptrn2_field = re.compile(r"\) .*$") + ptrn3_field = re.compile(r"\)\)$") + (field_name, n) = ptrn1_field.subn("", args, count=1) + if n != 0: + (field_name, n) = ptrn2_field.subn("", field_name, count=1) + if n == 0: + (field_name, n) = ptrn3_field.subn("", field_name, count=1) + return (field_name, n) + + +def set_field(args, field, val): + return re.subn( + field + r"\(\w+\)([ \)])", field + "(" + str(val) + ")\\1", args, count=1 + ) + + +hsa_patch_data = {} +ops_patch_data = {} + +# Fill API DB +api_table_descr = [ + [ + "BeginNs", + "EndNs", + "pid", + "tid", + "Name", + "args", + "Index", + "Data", + "__section", + "__lane", + ], + { + "BeginNs": "INTEGER", + "EndNs": "INTEGER", + "pid": "INTEGER", + "tid": "INTEGER", + "Name": "TEXT", + "args": "TEXT", + "Index": "INTEGER", + "Data": "TEXT", + "__section": "INTEGER", + "__lane": "INTEGER", + }, +] + + +# Filling API records DB table +# table_name - created DB table name +# db - DB handle +# indir - input directory +# api_name - traced API name +# api_pid - assigned JSON PID +# dep_pid - PID of dependet domain +# dep_list - list of dependet dospatch events +# dep_filtr - registered dependencies by record ID +def fill_api_db( + table_name, db, indir, api_name, api_pid, dep_pid, dep_list, dep_filtr, expl_id +): + global hsa_activity_found + global memory_manager + + range_start_times = {} + copy_csv = "" + + # Matches normal API records. + ptrn_api_record = re.compile(r"(\d+):(\d+) (\d+):(\d+) ([^\(]+)(\(.*)$") + # Matches records with a function name of "unknown" and no parameters. + # Capture groups 1-4 should match the same information as in ptrn_api_record. + # Used to avoid modifying ptrn_api_record regex. + ptrn_api_record_unknown = re.compile(r"(\d+):(\d+) (\d+):(\d+) (unknown).*$") + + hip_mcopy_ptrn = re.compile(r"hipMemcpy|hipMemset") + hip_wait_event_ptrn = re.compile(r"WaitEvent") + hip_sync_event_ptrn = re.compile(r"hipStreamSynchronize") + hip_sync_dev_event_ptrn = re.compile(r"hipDeviceSynchronize") + hip_graph_ptrn = re.compile(r"hipGraphLaunch") + wait_event_ptrn = re.compile(r"WaitEvent|hipStreamSynchronize|hipDeviceSynchronize") + hip_stream_wait_write_ptrn = re.compile( + r"hipStreamWaitValue64|hipStreamWriteValue64|hipStreamWaitValue32|hipStreamWriteValue32" + ) + prop_pattern = re.compile(r"([\w-]+)\((\w+)\)") + beg_pattern = re.compile(r'^dispatch\[(\d*)\], (.*) kernel-name\("([^"]*)"\)') + hip_strm_cr_event_ptrn = re.compile(r"hipStreamCreate") + hsa_mcopy_ptrn = re.compile(r"hsa_amd_memory_async_copy") + ptrn_fixformat = re.compile(r"(\d+:\d+ \d+:\d+ \w+)\(\s*(.*)\)$") + ptrn_fixkernel = re.compile(r"\s+kernel=(.*)$") + ptrn_multi_kernel = re.compile(r"(.*):(\d+)$") + ptrn_corr_id = re.compile(r"\ :(\d*)$") + + file_name = indir + "/" + api_name + "_api_trace.txt" + if not os.path.isfile(file_name): + return 0 + + hsa_copy_file_name = indir + "/" + "async_copy_trace.txt" + hsa_copy_file_name_present = 1 if os.path.isfile(file_name) else 0 + hsa_copy_deps = 1 if (api_pid == HSA_PID and hsa_copy_file_name_present == 1) else 0 + print("hsa_copy_deps: " + str(hsa_copy_deps)) + + # parsing an input trace file and creating a DB table + record_id_dict = {} + table_handle = db.add_table(table_name, api_table_descr) + with open(file_name, mode="r", errors="replace") as fd: + file_lines = fd.readlines() + total_lines = len(file_lines) + line_index = 0 + for line in file_lines: + if (line_index == total_lines - 1) or (line_index % 100 == 0): + sys.stdout.write( + "\rscan " + + api_name + + " API data " + + str(line_index) + + ":" + + str(total_lines) + + " " * 100 + ) + line_index += 1 + + record = line[:-1] + + corr_id = 0 + m = ptrn_corr_id.search(record) + if m: + corr_id = int(m.group(1)) + record = ptrn_corr_id.sub("", record) + + kernel_arg = "" + m = ptrn_fixkernel.search(record) + if m: + kernel_arg = "kernel(" + m.group(1) + ") " + record = ptrn_fixkernel.sub("", record) + + mfixformat = ptrn_fixformat.match(record) + if mfixformat: # replace '=' in args with parentheses + reformated_args = ( + kernel_arg + + mfixformat.group(2) + .replace("=", "(") + .replace(",", ")") + .replace("\\", "\\\\") + .replace('"', '\\"') + + ")" + ) + record = mfixformat.group(1) + "( " + reformated_args + ")" + + m = ptrn_api_record.match(record) + if not m: + m = ptrn_api_record_unknown.match(record) + if not m: + fatal(api_name + " bad record: '" + record + "'") + + rec_vals = [] + rec_len = len(api_table_descr[0]) - 3 + for ind in range(1, rec_len): + try: + rec_vals.append(m.group(ind)) + except IndexError: + rec_vals.append(None) + proc_id = int(rec_vals[2]) + thread_id = int(rec_vals[3]) + record_name = rec_vals[4] + # record_args is optional, and may be None if an unknown record is found. + record_args = rec_vals[5] + + # incrementing per-process record id/correlation id + if not proc_id in record_id_dict: + record_id_dict[proc_id] = 0 + record_id_dict[proc_id] += 1 + record_id = record_id_dict[proc_id] + + # setting correlationid to record id if correlation id is not defined + if corr_id == 0: + corr_id = record_id + + rec_vals.append(corr_id) + # extracting/converting stream id + (stream_id, stream_found) = get_field(record_args, "stream") + if stream_found: + stream_id = get_stream_index(stream_id) + (rec_vals[5], found) = set_field(record_args, "stream", stream_id) + if found == 0: + fatal( + 'set_field() failed for "stream", args: "' + record_args + '"' + ) + else: + (stream_id, stream_found) = get_field(record_args, "hStream") + if stream_found: + stream_id = get_stream_index(stream_id) + (rec_vals[5], found) = set_field( + record_args, "hStream", stream_id + ) + if found == 0: + fatal( + 'set_field() failed for "stream", args: "' + + record_args + + '"' + ) + else: + stream_id = 0 + + if hip_strm_cr_event_ptrn.match(record_name): + hip_streams.append(stream_id) + + if hip_sync_event_ptrn.match(record_name): + if (proc_id, stream_id) in last_hip_api_map: + (last_hip_api_corr_id, last_hip_api_from_pid) = last_hip_api_map[ + (proc_id, stream_id) + ][-1] + sync_api_beg_us = int((int(rec_vals[0]) - START_NS) / 1000) + if not proc_id in dep_dict: + dep_dict[proc_id] = {} + if HIP_PID not in dep_dict[proc_id]: + dep_dict[proc_id][HIP_PID] = { + "pid": last_hip_api_from_pid, + "from": [], + "to": {}, + "id": [], + } + dep_dict[proc_id][HIP_PID]["from"].append( + (-1, stream_id, thread_id) + ) + dep_dict[proc_id][HIP_PID]["id"].append(last_hip_api_corr_id) + dep_dict[proc_id][HIP_PID]["to"][ + last_hip_api_corr_id + ] = sync_api_beg_us + from_ids[(last_hip_api_corr_id, proc_id)] = ( + len(dep_dict[proc_id][HIP_PID]["from"]) - 1 + ) + + m = beg_pattern.match(record) + gpu_id = 0 + if m: + kernel_properties = m.group(2) + for prop in kernel_properties.split(", "): + m = prop_pattern.match(prop) + if m: + val = m.group(2) + var = m.group(1) + if var == "gpu-id": + gpu_id = int(val) + + if hsa_mcopy_ptrn.match(record_name) or hip_mcopy_ptrn.match(record_name): + ops_section_id = COPY_PID + else: + ops_section_id = GPU_BASE_PID + int(gpu_id) + + if (proc_id, stream_id) not in last_hip_api_map: + last_hip_api_map[(proc_id, stream_id)] = [] + last_hip_api_map[(proc_id, stream_id)].append((corr_id, ops_section_id)) + + # asyncronous opeartion API found + op_found = 0 + mcopy_found = 0 + + # extract kernel name string + (kernel_str, kernel_found) = get_field(record_args, "kernel") + if kernel_found == 0: + kernel_str = "" + else: + op_found = 1 + + if hip_mcopy_ptrn.match(record_name): + mcopy_found = 1 + op_found = 1 + + # HIP Graph API + if hip_graph_ptrn.search(record_name): + op_found = 1 + + # HIP WaitEvent API + if wait_event_ptrn.search(record_name): + op_found = 1 + + if hip_stream_wait_write_ptrn.search(record_name): + op_found = 1 + + # HSA memcopy API + if hsa_mcopy_ptrn.match(record_name): + mcopy_found = 1 + op_found = 1 + + stream_id = thread_id + hsa_patch_data[(corr_id, proc_id)] = thread_id + + if op_found: + roctx_msg = "" + + if not thread_id in range_start_times: + range_start_times[thread_id] = ( + sorted(range_data[thread_id].keys()) + if thread_id in range_data + else [] + ) + start_times = range_start_times[thread_id] + + index = bisect.bisect_right(start_times, int(rec_vals[0])) + if index > 0: + # We found the range that is closest to this operation. Iterate the + # range stack this range is part of until we find a range that entirely + # contains the operation. + range_start = start_times[index - 1] + while range_start != 0: + (range_end, range_start, msg) = range_data[thread_id][ + range_start + ] + if int(rec_vals[1]) < range_end: + # This range contains the operation. + roctx_msg = msg + break + + ops_patch_data[(corr_id, proc_id)] = ( + thread_id, + stream_id, + kernel_str, + roctx_msg, + ) + + if op_found: + op_found = 0 + beg_ns = int(rec_vals[0]) + end_ns = int(rec_vals[1]) + dur_us = int((end_ns - beg_ns) / 1000) + from_us = int((beg_ns - START_NS) / 1000) + dur_us / 2 + if api_pid == HIP_PID or hsa_copy_deps == 1: + if not proc_id in dep_dict: + dep_dict[proc_id] = {} + dep_proc = dep_dict[proc_id] + if not dep_pid in dep_proc: + if api_pid == "HIP_PID": + dep_proc[dep_pid] = {"pid": api_pid, "from": [], "id": []} + else: + dep_proc[dep_pid] = { + "pid": api_pid, + "from": [], + "id": [], + "to": {}, + } + dep_str = dep_proc[dep_pid] + dep_str["from"].append((from_us, stream_id, thread_id)) + if expl_id: + dep_str["id"].append(corr_id) + + # memcopy registering + api_data = ( + memory_manager.register_api(rec_vals) if mcopy_data_enabled else "" + ) + rec_vals.append(api_data) + + # setting section and lane + rec_vals.append(api_pid) # __section + rec_vals.append(thread_id) # __lane + + # inserting an API record to DB + db.insert_entry(table_handle, rec_vals) + + # inserting of dispatch events correlated to the dependent dispatches + for from_ns, proc_id, thread_id in dep_list: + if not proc_id in record_id_dict: + record_id_dict[proc_id] = 0 + record_id_dict[proc_id] += 1 + corr_id = record_id_dict[proc_id] + db.insert_entry( + table_handle, + [ + from_ns, + from_ns, + proc_id, + thread_id, + "hsa_dispatch", + "", + corr_id, + "", + api_pid, + thread_id, + ], + ) + + # generating memcopy CSV + if copy_csv != "": + file_name = os.environ["PWD"] + "/results_mcopy.csv" + with open(file_name, mode="w") as fd: + print("File '" + file_name + "' is generating") + fd.write(copy_csv) + + return 1 + + +############################################################# + +# fill COPY DB +copy_table_descr = [ + ["BeginNs", "EndNs", "Name", "pid", "tid", "Index", "Data", "__section", "__lane"], + { + "Index": "INTEGER", + "Name": "TEXT", + "args": "TEXT", + "BeginNs": "INTEGER", + "EndNs": "INTEGER", + "pid": "INTEGER", + "tid": "INTEGER", + "Data": "TEXT", + "__section": "INTEGER", + "__lane": "INTEGER", + }, +] + + +def fill_copy_db(table_name, db, indir): + sect_id = COPY_PID + file_name = indir + "/" + "async_copy_trace.txt" + ptrn_val = re.compile(r"^(\d+):(\d+) (async-copy):(\d+):(\d+)$") + + if not os.path.isfile(file_name): + return 0 + + table_handle = db.add_table(table_name, copy_table_descr) + with open(file_name, mode="r", errors="replace") as fd: + for line in fd.readlines(): + record = line[:-1] + m = ptrn_val.match(record) + if not m: + fatal("bad async-copy entry '" + record + "'") + else: + rec_vals = [] + for ind in range(1, 4): + rec_vals.append(m.group(ind)) + corr_id = int(m.group(4)) + proc_id = int(m.group(5)) + + # querying tid value + if (corr_id, proc_id) in hsa_patch_data: + thread_id = hsa_patch_data[(corr_id, proc_id)] + else: + thread_id = -1 + + # completing record + rec_vals.append(proc_id) # tid + rec_vals.append(thread_id) # tid + rec_vals.append(corr_id) # Index + + # registering memcopy information + activity_data = ( + memory_manager.register_copy(rec_vals) if mcopy_data_enabled else "" + ) + rec_vals.append(activity_data) + + # appending straem ID and section ID + rec_vals.append(COPY_PID) # __section + rec_vals.append(thread_id) # __lane + + # inserting DB activity entry + db.insert_entry(table_handle, rec_vals) + + # filling dependencies + to_ns = int(rec_vals[0]) + to_us = int((to_ns - START_NS) / 1000) + + if thread_id != -1: + # if not proc_id in dep_dict: dep_dict[proc_id] = {} + dep_proc = dep_dict[proc_id] + # if not pid in dep_proc: dep_proc[pid] = { 'pid': HSA_PID, 'from': [], 'to': {}, 'id': [] } + dep_str = dep_proc[sect_id] + dep_str["to"][corr_id] = to_us + dep_str["id"].append(corr_id) + + return 1 + + +############################################################# + +# fill HCC ops DB +ops_table_descr = [ + [ + "BeginNs", + "EndNs", + "dev-id", + "queue-id", + "Name", + "pid", + "tid", + "roctx-range", + "stream-id", + "Index", + "Data", + "__section", + "__lane", + ], + { + "Index": "INTEGER", + "Name": "TEXT", + "args": "TEXT", + "BeginNs": "INTEGER", + "EndNs": "INTEGER", + "dev-id": "INTEGER", + "queue-id": "INTEGER", + "pid": "INTEGER", + "tid": "INTEGER", + "roctx-range": "TEXT", + "Data": "TEXT", + "stream-id": "INTEGER", + "__section": "INTEGER", + "__lane": "INTEGER", + }, +] + + +def fill_ops_db(kernel_table_name, mcopy_table_name, db, indir): + global max_gpu_id + file_name = indir + "/" + "hcc_ops_trace.txt" + ptrn_val = re.compile(r"(\d+):(\d+) (\d+):(\d+) (.*)$") + ptrn_id = re.compile(r"^([^:]+):(\d+):(\d+)$") + ptrn_mcopy = re.compile(r"(Memcpy|Copy|Fill)") + ptrn_barrier = re.compile(r"Marker") + + if not os.path.isfile(file_name): + return {} + + filtr = {} + + kernel_table_handle = db.add_table(kernel_table_name, ops_table_descr) + mcopy_table_handle = db.add_table(mcopy_table_name, ops_table_descr) + with open(file_name, mode="r", errors="replace") as fd: + file_lines = fd.readlines() + total_lines = len(file_lines) + line_index = 0 + for line in file_lines: + if (line_index == total_lines - 1) or (line_index % 100 == 0): + sys.stdout.write( + "\rscan ops data " + + str(line_index) + + ":" + + str(total_lines) + + " " * 100 + ) + line_index += 1 + + record = line[:-1] + m = ptrn_val.match(record) + if m: + # parsing trace record + rec_vals = [] + for ind in range(1, 6): + rec_vals.append(m.group(ind)) + label = rec_vals[4] # record name + m = ptrn_id.match(label) + if not m: + fatal("bad hcc ops entry '" + record + "'") + name = m.group(1) + corr_id = int(m.group(2)) + proc_id = int(m.group(3)) + + # checking name for memcopy pattern + is_barrier = 0 + if ptrn_mcopy.search(name): + rec_table_name = mcopy_table_name + table_handle = mcopy_table_handle + sect_id = COPY_PID + else: + rec_table_name = kernel_table_name + table_handle = kernel_table_handle + + gpu_id = int(rec_vals[2]) + if gpu_id > max_gpu_id: + max_gpu_id = gpu_id + + if ptrn_barrier.search(name): + name = '""' + is_barrier = 1 + sect_id = GPU_BASE_PID + int(gpu_id) + 512 + else: + sect_id = GPU_BASE_PID + int(gpu_id) + + thread_id = 0 + stream_id = 0 + roctx_range = "" + if (corr_id, proc_id) in ops_patch_data: + (thread_id, stream_id, name_patch, roctx_range) = ops_patch_data[ + (corr_id, proc_id) + ] + if name_patch != "": + name = name_patch + if roctx_range == "": + roctx_range = name + else: + if is_barrier: + continue + else: + if "ROCP_CTRL_RATE" in os.environ: + continue + else: + fatal( + "hcc ops data not found: '" + + record + + "', " + + str(corr_id) + + ", " + + str(proc_id) + ) + + # activity record + rec_vals[4] = name # Name + rec_vals.append(proc_id) # pid + rec_vals.append(thread_id) # tid + rec_vals.append(roctx_range) # roctx-range + rec_vals.append(stream_id) # StreamId + rec_vals.append(corr_id) # Index + + # registering memcopy information + activity_data = ( + memory_manager.register_activity(rec_vals) + if mcopy_data_enabled + else "" + ) + rec_vals.append(activity_data) + + # activity record data for stream ID and sction ID + rec_vals.append(sect_id) # __section + rec_vals.append(stream_id) # __lane + + # inserting DB activity entry + db.insert_entry(table_handle, rec_vals) + + # registering a dependency filtr + filtr[(corr_id, proc_id)] = rec_table_name + + # filling a dependencies + to_ns = int(rec_vals[0]) + to_us = int((to_ns - START_NS) / 1000) + + end_ns = int(rec_vals[1]) + dur_us = int((end_ns - to_ns) / 1000) + + if (corr_id, proc_id) in from_ids: + depid = from_ids[(corr_id, proc_id)] + from_val = dep_dict[proc_id][HIP_PID]["from"][depid] + print("from_val" + str(from_val)) + from_val_new = (to_us + dur_us, from_val[1], from_val[2]) + dep_dict[proc_id][HIP_PID]["from"][depid] = from_val_new + + if not proc_id in dep_dict: + dep_dict[proc_id] = {} + dep_proc = dep_dict[proc_id] + if not sect_id in dep_proc: + dep_proc[sect_id] = {"bsp": OPS_PID, "to": {}} + dep_str = dep_proc[sect_id] + dep_str["to"][corr_id] = to_us + + else: + fatal("hcc ops bad record: '" + record + "'") + + return filtr + + +############################################################# +# main +if len(sys.argv) < 2: + fatal("Usage: " + sys.argv[0] + " ") + +outfile = sys.argv[1] +infiles = sys.argv[2:] +indir = re.sub(r"\/[^\/]*$", r"", infiles[0]) +inext = re.sub(r"\s+$", r"", infiles[0]) +inext = re.sub(r"^.*(\.[^\.]+)$", r"\1", inext) + +dbfile = "" +csvfile = "" + +if "ROCP_JSON_REBASE" in os.environ and os.environ["ROCP_JSON_REBASE"] == 0: + begin_ts_file = indir + "/begin_ts_file.txt" + if os.path.isfile(begin_ts_file): + with open(begin_ts_file, mode="r", errors="replace") as fd: + ind = 0 + for line in fd.readlines(): + val = int(line) + if ind == 0 or val < START_NS: + START_NS = val + ind += 1 + print("START timestamp found (" + str(START_NS) + "ns)") + +if re.search(r"\.csv$", outfile): + csvfile = outfile +elif re.search(r"\.db$", outfile): + dbfile = outfile + csvfile = re.sub(r"\.db$", ".csv", outfile) +else: + fatal("Bad output file '" + outfile + "'") + +if inext == ".txt": + for f in infiles: + parse_res(f) + if len(var_table) != 0: + merge_table() + +if dbfile == "": + dump_csv(csvfile) +else: + statfile = re.sub(r"\.csv$", ".stats.csv", csvfile) + jsonfile = re.sub(r"\.csv$", ".json", csvfile) + + hsa_statfile = re.sub(r"\.stats\.csv$", r".hsa_stats.csv", statfile) + hip_statfile = re.sub(r"\.stats\.csv$", r".hip_stats.csv", statfile) + ops_statfile = statfile + copy_statfile = re.sub(r"\.stats\.csv$", r".copy_stats.csv", statfile) + memcopy_info_file = re.sub(r"\.stats\.csv$", r".memcopy_info.csv", statfile) + sysinfo_file = re.sub(r"\.stats\.csv$", r".sysinfo.txt", statfile) + metadata_gen(sysinfo_file, "rocminfo") + + with open(dbfile, mode="w") as fd: + fd.truncate() + db = SQLiteDB(dbfile) + memory_manager = MemManager(db, indir) + + ext_trace_found = fill_ext_db("rocTX", db, indir, "roctx", EXT_PID) + + hsa_trace_found = fill_api_db( + "HSA", db, indir, "hsa", HSA_PID, COPY_PID, kern_dep_list, {}, 0 + ) + hsa_activity_found = fill_copy_db("COPY", db, indir) + + hip_trace_found = fill_api_db("HIP", db, indir, "hip", HIP_PID, OPS_PID, [], {}, 1) + ops_filtr = fill_ops_db("OPS", "COPY", db, indir) + + fill_kernel_db("KERN", db) + + any_trace_found = ext_trace_found | hsa_trace_found | hip_trace_found + copy_trace_found = 0 + if hsa_activity_found or len(ops_filtr): + copy_trace_found = 1 + + db.open_json(jsonfile) + + if ext_trace_found: + db.label_json(EXT_PID, "Markers and Ranges", jsonfile) + + if hip_trace_found: + db.label_json(HIP_PID, "CPU HIP API", jsonfile) + + if hsa_trace_found: + db.label_json(HSA_PID, "CPU HSA API", jsonfile) + + db.label_json(COPY_PID, "COPY", jsonfile) + + if max_gpu_id >= 0: + for ind in range(0, int(max_gpu_id) + 1): + db.label_json(int(ind) + int(GPU_BASE_PID), "GPU" + str(ind), jsonfile) + db.label_json(int(ind) + int(GPU_BASE_PID) + 512 , "GPU Barriers" + str(ind), jsonfile) + + if ext_trace_found: + dform.gen_ext_json_trace(db, "rocTX", START_NS, jsonfile) + + if len(var_table) != 0: + dform.post_process_data(db, "KERN", csvfile) + dform.gen_table_bins(db, "KERN", statfile, "KernelName", "DurationNs") + if "BeginNs" in var_list: + dform.gen_kernel_json_trace(db, "KERN", GPU_BASE_PID, START_NS, jsonfile) + + if hsa_trace_found: + dform.post_process_data(db, "HSA") + dform.gen_table_bins(db, "HSA", hsa_statfile, "Name", "DurationNs") + dform.gen_api_json_trace(db, "HSA", START_NS, jsonfile) + + if copy_trace_found: + dform.post_process_data(db, "COPY") + dform.gen_table_bins(db, "COPY", copy_statfile, "Name", "DurationNs") + dform.gen_api_json_trace(db, "COPY", START_NS, jsonfile) + + if hip_trace_found: + dform.post_process_data(db, "HIP") + dform.gen_table_bins(db, "HIP", hip_statfile, "Name", "DurationNs") + dform.gen_api_json_trace(db, "HIP", START_NS, jsonfile) + + if ops_filtr: + dform.post_process_data(db, "OPS") + dform.gen_table_bins(db, "OPS", ops_statfile, "Name", "DurationNs") + dform.gen_ops_json_trace(db, "OPS", GPU_BASE_PID, START_NS, jsonfile) + + if any_trace_found: + dep_id = 0 + for proc_id, dep_proc in dep_dict.items(): + for to_pid, dep_str in dep_proc.items(): + if "bsp" in dep_str: + bspid = dep_str["bsp"] + base_str = dep_proc[bspid] + for v in ("pid", "from", "id"): + dep_str[v] = base_str[v] + base_str["inv"] = 1 + + for to_pid, dep_str in dep_proc.items(): + if "inv" in dep_str: + continue + if not "to" in dep_str: + continue + + from_pid = dep_str["pid"] + from_us_list = dep_str["from"] + to_us_dict = dep_str["to"] + corr_id_list = dep_str["id"] + + db.flow_json( + dep_id, + from_pid, + from_us_list, + to_pid, + to_us_dict, + corr_id_list, + jsonfile, + ) + dep_id += len(from_us_list) + + db.metadata_json(jsonfile, sysinfo_file) + db.close_json(jsonfile) + + if mcopy_data_enabled: + memory_manager.dump_data("MM", memcopy_info_file) + + db.close() + +sys.exit(0) +############################################################# diff --git a/projects/rocprofiler/bin/txt2params.py b/projects/rocprofiler/bin/txt2params.py new file mode 100644 index 0000000000..4be3426653 --- /dev/null +++ b/projects/rocprofiler/bin/txt2params.py @@ -0,0 +1,113 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +import os, sys, re + + +# gen_params() takes a text file like the output of rocminfo cmd and parses it into a map {key,value} +# where key is the param and value is the value of this param +# for example: Threadmodel : "posix" +# it also processes encompasing sections to generate a full param name such as (section names separated by '_'): +# "Agent2_PoolInfo_ISAInfo_ISA1_WorkgroupMaxSizeperDimension_x": "1024(0x400)", +def gen_params(txtfile): + fields = {} + counter = 0 + parent_field = "" + nbr_indent = 0 + nbr_indent_prev = 0 + check_for_dims = False + with open(txtfile) as fp: + for line in fp: + me = re.match(r"\*\*\* Done \*\*\*", line) # Marks the end of cmd + if me: + parent_field = "" + nbr_indent = 0 + nbr_indent_prev = 0 + check_for_dims = False + continue + mv = re.match( + r"HCC clang version\s+(.*)", line + ) # outlier: only line with a version number and no ':', special case + if mv: + key = "HCCclangversion" + val = mv.group(1) + counter = counter + 1 + fields[(counter, key)] = val + continue + # Variable 'check_for_dims' is True for text like this: + # Workgroup Max Size per Dimension: + # x 1024(0x400) + # y 1024(0x400) + # z 1024(0x400) + if check_for_dims == True: + mc = re.match(r"\s*([x|y|z])\s+(.*)", line) + if mc: + key_sav = mc.group(1) + if parent_field != "": + key = parent_field + "." + mc.group(1) + else: + key = mc.group(1) + val = re.sub(r"\s+", "", mc.group(2)) + counter = counter + 1 + fields[(counter, key)] = val + if key_sav == "z": + check_for_dims = False + nbr_indent_prev = nbr_indent + mi = re.search(r"^(\s+)\w+.*", line) + md = re.search(":", line) + if mi: + nbr_indent = int(len(mi.group(1)) / 2) # indentation cnt + else: + if not md: + tmp = re.sub(r"\s+", "", line) + if tmp.isalnum(): + parent_field = tmp + + if nbr_indent < nbr_indent_prev: + go_back_parent = nbr_indent_prev - nbr_indent + for i in range(go_back_parent): # decrease as many levels up as needed + pos = parent_field.rfind(".") + if pos != -1: + parent_field = parent_field[:pos] + # Process lines such as : + # Segment: GLOBAL; FLAGS: KERNARG, FINE GRAINED + # Size: 131897644(0x7dc992c) KB + for lin in line.split(";"): + lin = re.sub(r"\s+", "", lin) + m = re.match(r"(.*):(.*)", lin) + if m: + key, val = m.group(1), m.group(2) + if parent_field != "": + key = parent_field + "." + key + if val == "": + mk = re.match(r".*Dimension", key) + if mk: # expect x,y,z on next 3 lines + check_for_dims = True + parent_field = key + else: + counter = counter + 1 + fields[(counter, key)] = val + else: + if nbr_indent != nbr_indent_prev and not check_for_dims: + parent_field = parent_field + "." + lin.replace(":", "") + + return fields diff --git a/projects/rocprofiler/bin/txt2xml.sh b/projects/rocprofiler/bin/txt2xml.sh new file mode 100755 index 0000000000..c090ac4380 --- /dev/null +++ b/projects/rocprofiler/bin/txt2xml.sh @@ -0,0 +1,103 @@ +#!/bin/bash + +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +timestamp=`date +%y%m%d_%H%M%S` + +if [ $# = 0 ] ; then + echo "Usage: $0 [output dir]" + exit -1 +fi + +input=$1 +outdir=$2 +if [ -z "$outdir" ] ; then + outdir="." +fi + +range="" +kernel="" +gpu_index="" + +parse() { + scan="$1" + index=0 + while read -r line || [[ -n "$line" ]] ; do + line=`echo $line | sed "s/\s*#.*$//"` + if [ -z "$line" ] ; then + continue + fi + + feature=`echo $line | sed -n "s/^\s*\([a-z]*\)\s*:.*$/\1/p"` + line=`echo $line | sed "s/^[^:]*:\s*//"` + line=`echo "$line" | sed -e "s/\s*=\s*/=/g" -e "s/\s*:\s*/:/g" -e "s/,\{1,\}/ /g" -e "s/\s\{1,\}/ /g" -e "s/\s*$//"` + + if [ "$scan" = 0 ] ; then + line=`echo "$line" | sed -e "s/ /,/g"` + if [ "$feature" == "range" ] ; then + range=$line + fi + if [ "$feature" == "kernel" ] ; then + kernel=$line + fi + if [ "$feature" == "gpu" ] ; then + gpu_index=$line + fi + else + found=$(echo $feature | sed -n "/^\(pmc\|hip\|hsa\)$/ p") + if [ -n "$found" ] ; then + output=$outdir/input${index}.xml + header="# $timestamp '$output' generated with '$0 $*'" + echo $header > $output + + if [ "$feature" == "pmc" ] ; then + line=`echo "$line" | sed -e "s/ /,/g"` + cat >> $output < + +EOF + fi + + if [ "$feature" == "hip" ] ; then + cat >> $output < +EOF + fi + + if [ "$feature" == "hsa" ] ; then + cat >> $output < +EOF + fi + + fi + fi + + index=$((index + 1)) + done < $input +} + +parse 0 +parse 1 + +exit 0 diff --git a/projects/rocprofiler/build.sh b/projects/rocprofiler/build.sh new file mode 100755 index 0000000000..64c776a362 --- /dev/null +++ b/projects/rocprofiler/build.sh @@ -0,0 +1,109 @@ +#!/bin/bash -e + +################################################################################ +# Copyright (c) 2018-2022 Advanced Micro Devices, Inc. +# +# 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. +################################################################################ + +SRC_DIR=$(dirname "$0") +COMPONENT="rocprofiler" +ROCM_PATH="${ROCM_PATH:=/opt/rocm}" +LD_RUNPATH_FLAG=" -Wl,--enable-new-dtags -Wl,--rpath,$ROCM_PATH/lib" + +usage() { + echo -e "ROCProfiler Build Script Usage:" + echo -e "\nTo run ./run.sh PARAMs, PARAMs can be the following:" + echo -e "-h | --help For showing this message" + echo -e "-b | --build For compiling" + echo -e "-cb | --clean-build For full clean build" + echo -e "-acb | --asan-clean-build For compiling with ASAN library attached" + exit 1 +} + +while [ 1 ] ; do + if [[ "$1" = "-h" || "$1" = "--help" ]] ; then + usage + exit 1 + elif [[ "$1" = "-b" || "$1" = "--build" ]] ; then + TO_CLEAN=no + shift + elif [[ "$1" = "-acb" || "$1" = "--asan-clean-build" ]] ; then + ASAN=True TO_CLEAN=yes + shift + elif [[ "$1" = "-cb" || "$1" = "--clean-build" ]] ; then + TO_CLEAN=yes + shift + elif [ "$1" = "--" ] ; then + shift + EXTRA_BUILD_ARGS=$@ + break + elif [[ "$1" = "-"* ]] ; then + echo -e "Wrong option \"$1\", Please use the following options:\n" + usage + exit 1 + else + break + fi +done + +umask 022 + +if [ -z "$ROCPROFILER_ROOT" ]; then ROCPROFILER_ROOT=$SRC_DIR; fi +if [ -z "$BUILD_DIR" ] ; then BUILD_DIR=build; fi +if [ -z "$BUILD_TYPE" ] ; then BUILD_TYPE="RelWithDebInfo"; fi +if [ -z "$PACKAGE_ROOT" ] ; then PACKAGE_ROOT=$ROCM_PATH; fi +if [ -z "$PREFIX_PATH" ] ; then PREFIX_PATH=$PACKAGE_ROOT; fi +if [ -z "$HIP_VDI" ] ; then HIP_VDI=0; fi +if [ -n "$ROCM_RPATH" ] ; then LD_RUNPATH_FLAG=" -Wl,--enable-new-dtags -Wl,--rpath,${ROCM_RPATH}"; fi +if [ -z "$TO_CLEAN" ] ; then TO_CLEAN=yes; fi +if [ -z "$RUN_TEST" ] ; then RUN_TEST=no; fi +if [ -z "$ASAN" ] ; then ASAN=False; fi +if [ -z "$GPU_LIST" ] ; then GPU_LIST="gfx900 gfx906 gfx908 gfx90a gfx940 gfx941 gfx942 gfx1030 gfx1100 gfx1101 gfx1102 gfx1031 gfx1150 gfx1151 gfx1200 gfx1201"; fi + +ROCPROFILER_ROOT=$(cd $ROCPROFILER_ROOT && echo $PWD) + +if [ "$TO_CLEAN" = "yes" ] ; then rm -rf $BUILD_DIR; fi + +cmake -B ${BUILD_DIR} ${ROCPROFILER_ROOT} \ + -DROCPROFILER_BUILD_CI=1 \ + -DROCPROFILER_BUILD_TESTS=1 \ + -DROCPROFILER_BUILD_SAMPLES=1 \ + -DCMAKE_EXPORT_COMPILE_COMMANDS=TRUE \ + -DCMAKE_BUILD_TYPE=${BUILD_TYPE:-'RelWithDebInfo'} \ + -DCMAKE_MODULE_PATH="${ROCM_PATH}/hip/cmake;${ROCM_PATH}/lib/cmake" \ + -DCMAKE_PREFIX_PATH="${ROCM_INSTALL_PATH}/llvm;$PREFIX_PATH" \ + -DCMAKE_INSTALL_PREFIX="$PACKAGE_ROOT" \ + -DCMAKE_SHARED_LINKER_FLAGS="$LD_RUNPATH_FLAG" \ + -DCMAKE_INSTALL_RPATH=${ROCM_RPATH} \ + -DCMAKE_INSTALL_LIBDIR=lib \ + -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=FALSE \ + -DGPU_TARGETS="$GPU_LIST" \ + -DCPACK_PACKAGING_INSTALL_PREFIX=$PACKAGE_ROOT \ + -DCPACK_GENERATOR=${CPACKGEN:-'DEB;RPM'} \ + -DCPACK_OBJCOPY_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-objcopy" \ + -DCPACK_READELF_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-readelf" \ + -DCPACK_STRIP_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-strip" \ + -DCPACK_OBJDUMP_EXECUTABLE="${PACKAGE_ROOT}/llvm/bin/llvm-objdump" \ + ${EXTRA_BUILD_ARGS} + +cmake --build "$BUILD_DIR" --target all --parallel $(nproc) +cmake --build "$BUILD_DIR" --target package --parallel $(nproc) + +exit 0 diff --git a/projects/rocprofiler/cmake_modules/FindLibDw.cmake b/projects/rocprofiler/cmake_modules/FindLibDw.cmake new file mode 100644 index 0000000000..8793fb8196 --- /dev/null +++ b/projects/rocprofiler/cmake_modules/FindLibDw.cmake @@ -0,0 +1,21 @@ +# Try to find LIBDW +# +# Once found, this will define: +# - LIBDW_FOUND - system has libelf +# - LIBDW_INCLUDE_DIRS - the libelf include directory +# - LIBDW_LIBRARIES - Link these to use libelf +# - LIBDW_DEFINITIONS - Compiler switches required for using libelf +find_path( + FIND_LIBDW_INCLUDES + NAMES elfutils/libdw.h + PATHS /usr/include /usr/local/include) + +find_library(FIND_LIBDW_LIBRARIES NAMES dw PATH /usr/lib /usr/local/lib) + +include(FindPackageHandleStandardArgs) +find_package_handle_standard_args(LibDw DEFAULT_MSG FIND_LIBDW_INCLUDES + FIND_LIBDW_LIBRARIES) +mark_as_advanced(FIND_LIBDW_INCLUDES FIND_LIBDW_LIBRARIES) + +set(LIBDW_INCLUDES ${FIND_LIBDW_INCLUDES}) +set(LIBDW_LIBRARIES ${FIND_LIBDW_LIBRARIES}) diff --git a/projects/rocprofiler/cmake_modules/FindLibElf.cmake b/projects/rocprofiler/cmake_modules/FindLibElf.cmake new file mode 100644 index 0000000000..6355b09ba2 --- /dev/null +++ b/projects/rocprofiler/cmake_modules/FindLibElf.cmake @@ -0,0 +1,21 @@ +# Try to find LIBELF +# +# Once found, this will define: +# - LIBELF_FOUND - system has libelf +# - LIBELF_INCLUDE_DIRS - the libelf include directory +# - LIBELF_LIBRARIES - Link these to use libelf +# - LIBELF_DEFINITIONS - Compiler switches required for using libelf +find_path( + FIND_LIBELF_INCLUDES + NAMES libelf.h + PATHS /usr/include /usr/include/libelf /usr/local/include /usr/local/include/libelf) + +find_library(FIND_LIBELF_LIBRARIES NAMES elf PATH /usr/lib /usr/local/lib) + +include(FindPackageHandleStandardArgs) +find_package_handle_standard_args(LibElf DEFAULT_MSG FIND_LIBELF_INCLUDES + FIND_LIBELF_LIBRARIES) +mark_as_advanced(FIND_LIBELF_INCLUDES FIND_LIBELF_LIBRARIES) + +set(LIBELF_INCLUDES ${FIND_LIBELF_INCLUDES}) +set(LIBELF_LIBRARIES ${FIND_LIBELF_LIBRARIES}) diff --git a/projects/rocprofiler/cmake_modules/Templates/gtest_discover_tests_properties.cmake.in b/projects/rocprofiler/cmake_modules/Templates/gtest_discover_tests_properties.cmake.in new file mode 100644 index 0000000000..14a751d972 --- /dev/null +++ b/projects/rocprofiler/cmake_modules/Templates/gtest_discover_tests_properties.cmake.in @@ -0,0 +1,11 @@ +cmake_minimum_required(VERSION 3.18.0 FATAL_ERROR) + +if(NOT @GTEST_DISCOVER_TESTS_TARGET@_TESTS) + message(FATAL_ERROR "@GTEST_DISCOVER_TESTS_TARGET@_TESTS is not defined") +endif() + +foreach(_TEST ${@GTEST_DISCOVER_TESTS_TARGET@_TESTS}) + set_tests_properties( + ${_TEST} PROPERTIES LABELS "@GTEST_DISCOVER_TESTS_LABELS@" ENVIRONMENT + "@GTEST_DISCOVER_TESTS_ENVIRONMENT@") +endforeach() diff --git a/projects/rocprofiler/cmake_modules/rocprofiler_env.cmake b/projects/rocprofiler/cmake_modules/rocprofiler_env.cmake new file mode 100644 index 0000000000..517b033238 --- /dev/null +++ b/projects/rocprofiler/cmake_modules/rocprofiler_env.cmake @@ -0,0 +1,76 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +add_library(rocprofiler-build-flags INTERFACE) +add_library(rocprofiler::build-flags ALIAS rocprofiler-build-flags) + +target_compile_options( + rocprofiler-build-flags + INTERFACE $<$:-W -Wall -Wextra -Wno-unused-parameter> + $<$:-fms-extensions> + $<$:$<$:-ferror-limit=1000000>> + ) +target_compile_definitions(rocprofiler-build-flags INTERFACE NEW_TRACE_API=1) + +# Enable debug trace +if(ROCPROFILER_DEBUG_TRACE) + target_compile_definitions(rocprofiler-build-flags INTERFACE DEBUG_TRACE=1) +endif() + +# Enable direct loading of AQL-profile HSA extension +if(ROCPROFILER_LD_AQLPROFILE) + target_compile_definitions(rocprofiler-build-flags INTERFACE ROCP_LD_AQLPROFILE=1) +endif() + +if(NOT DEFINED ROCM_PATH) + set(ROCM_PATH "/opt/rocm" CACHE STRING "Default ROCM installation directory") +endif() + +# Find hsa-runtime +find_package( + hsa-runtime64 CONFIG REQUIRED + HINTS ${CMAKE_PREFIX_PATH} + PATHS ${ROCM_PATH} + PATH_SUFFIXES lib/cmake/hsa-runtime64) + +# Include path: /opt/rocm-ver/include. Go up one level to get ROCm path +get_filename_component(ROCM_ROOT_DIR ${ROCM_PATH}/include DIRECTORY) + +# Basic Tool Chain Information +message("----------Build-Type: ${CMAKE_BUILD_TYPE}") +message("------------Compiler: ${CMAKE_CXX_COMPILER}") +message("----Compiler-Version: ${CMAKE_CXX_COMPILER_VERSION}") +message("-------ROCM_ROOT_DIR: ${ROCM_ROOT_DIR}") +message("-----CMAKE_CXX_FLAGS: ${CMAKE_CXX_FLAGS}") +message("---CMAKE_PREFIX_PATH: ${CMAKE_PREFIX_PATH}") +message("---------GPU_TARGETS: ${GPU_TARGETS}") + +if("${ROCM_ROOT_DIR}" STREQUAL "") + message(FATAL_ERROR "ROCM_ROOT_DIR is not found.") +endif() + +find_library( + HSA_AMD_AQLPROFILE_LIBRARY + NAMES hsa-amd-aqlprofile64 + HINTS ${CMAKE_PREFIX_PATH} + PATHS ${ROCM_ROOT_DIR} + PATH_SUFFIXES lib REQUIRED) diff --git a/projects/rocprofiler/cmake_modules/rocprofiler_formatting.cmake b/projects/rocprofiler/cmake_modules/rocprofiler_formatting.cmake new file mode 100644 index 0000000000..35192a3906 --- /dev/null +++ b/projects/rocprofiler/cmake_modules/rocprofiler_formatting.cmake @@ -0,0 +1,103 @@ +# ------------------------------------------------------------------------------# +# +# creates following targets to format code: +# - format +# - format-source +# - format-cmake +# - format-python +# - format-rocprofiler-source +# - format-rocprofiler-cmake +# - format-rocprofiler-python +# +# ------------------------------------------------------------------------------# + +include_guard(GLOBAL) + +find_program(ROCPROFILER_CLANG_FORMAT_EXE NAMES clang-format-11 clang-format-mp-11) +find_program(ROCPROFILER_CMAKE_FORMAT_EXE NAMES cmake-format) +find_program(ROCPROFILER_BLACK_FORMAT_EXE NAMES black) + +if(ROCPROFILER_CLANG_FORMAT_EXE + OR ROCPROFILER_BLACK_FORMAT_EXE + OR ROCPROFILER_CMAKE_FORMAT_EXE) + add_custom_target(format-rocprofiler) + + if(NOT TARGET format) + add_custom_target(format) + endif() + + foreach(_TYPE source python cmake) + if(NOT TARGET format-${_TYPE}) + add_custom_target(format-${_TYPE}) + endif() + endforeach() + + set(rocp_sources) + set(rocp_headers) + set(rocp_cmake_files) + set(rocp_python_files) + foreach(_DIR include src plugin samples test tests-v2 script cmake_modules) + foreach(_TYPE headers sources cmake_files python_files) + set(${_TYPE}) + endforeach() + file(GLOB_RECURSE headers ${PROJECT_SOURCE_DIR}/${_DIR}/*.h) + file(GLOB_RECURSE sources ${PROJECT_SOURCE_DIR}/${_DIR}/*.cpp) + file(GLOB_RECURSE cmake_files ${PROJECT_SOURCE_DIR}/${_DIR}/*CMakeLists.txt + ${PROJECT_SOURCE_DIR}/${_DIR}/*.cmake) + file(GLOB_RECURSE python_files ${PROJECT_SOURCE_DIR}/${_DIR}/*.py) + foreach(_TYPE headers sources cmake_files python_files) + list(APPEND rocp_${_TYPE} ${${_TYPE}}) + endforeach() + endforeach() + + if(ROCPROFILER_CLANG_FORMAT_EXE) + add_custom_target( + format-rocprofiler-source + ${ROCPROFILER_CLANG_FORMAT_EXE} -i ${rocp_sources} ${rocp_headers} + COMMENT + "[rocprofiler] Running source formatter ${ROCPROFILER_CLANG_FORMAT_EXE}..." + ) + endif() + + if(ROCPROFILER_BLACK_FORMAT_EXE) + add_custom_target( + format-rocprofiler-python + ${ROCPROFILER_BLACK_FORMAT_EXE} -q ${rocp_python_files} + COMMENT + "[rocprofiler] Running Python formatter ${ROCPROFILER_BLACK_FORMAT_EXE}..." + ) + if(NOT TARGET format-python) + add_custom_target(format-python) + endif() + endif() + + if(ROCPROFILER_CMAKE_FORMAT_EXE) + add_custom_target( + format-rocprofiler-cmake + ${ROCPROFILER_CMAKE_FORMAT_EXE} -i ${rocp_cmake_files} + COMMENT + "[rocprofiler] Running CMake formatter ${ROCPROFILER_CMAKE_FORMAT_EXE}..." + ) + if(NOT TARGET format-cmake) + add_custom_target(format-cmake) + endif() + endif() + + foreach(_TYPE source python cmake) + if(TARGET format-rocprofiler-${_TYPE}) + add_dependencies(format-rocprofiler format-rocprofiler-${_TYPE}) + add_dependencies(format-${_TYPE} format-rocprofiler-${_TYPE}) + endif() + endforeach() + + foreach(_TYPE source python) + if(TARGET format-rocprofiler-${_TYPE}) + add_dependencies(format format-rocprofiler-${_TYPE}) + endif() + endforeach() +else() + message( + STATUS + "no formatting tools (clang-format-11/black/cmake-format) could not be found. formatting build targets not available." + ) +endif() diff --git a/projects/rocprofiler/cmake_modules/rocprofiler_linting.cmake b/projects/rocprofiler/cmake_modules/rocprofiler_linting.cmake new file mode 100644 index 0000000000..3e00aff934 --- /dev/null +++ b/projects/rocprofiler/cmake_modules/rocprofiler_linting.cmake @@ -0,0 +1,30 @@ +include_guard(DIRECTORY) + +# ----------------------------------------------------------------------------------------# +# +# Clang Tidy +# +# ----------------------------------------------------------------------------------------# + +if(ROCPROFILER_ENABLE_CLANG_TIDY) + find_program(ROCPROFILER_CLANG_TIDY_COMMAND NAMES clang-tidy) + + if(NOT ROCPROFILER_CLANG_TIDY_COMMAND) + message( + WARNING "ROCPROFILER_ENABLE_CLANG_TIDY is ON but clang-tidy is not found!") + set(ROCPROFILER_ENABLE_CLANG_TIDY OFF) + else() + set(CMAKE_CXX_CLANG_TIDY ${ROCPROFILER_CLANG_TIDY_COMMAND} + -header-filter=${PROJECT_SOURCE_DIR}/.*) + + # Create a preprocessor definition that depends on .clang-tidy content so the + # compile command will change when .clang-tidy changes. This ensures that a + # subsequent build re-runs clang-tidy on all sources even if they do not otherwise + # need to be recompiled. Nothing actually uses this definition. We add it to + # targets on which we run clang-tidy just to get the build dependency on the + # .clang-tidy file. + file(SHA1 ${PROJECT_SOURCE_DIR}/.clang-tidy clang_tidy_sha1) + set(CLANG_TIDY_DEFINITIONS "CLANG_TIDY_SHA1=${clang_tidy_sha1}") + unset(clang_tidy_sha1) + endif() +endif() diff --git a/projects/rocprofiler/cmake_modules/rocprofiler_options.cmake b/projects/rocprofiler/cmake_modules/rocprofiler_options.cmake new file mode 100644 index 0000000000..41135b5aea --- /dev/null +++ b/projects/rocprofiler/cmake_modules/rocprofiler_options.cmake @@ -0,0 +1,139 @@ +if("${PROJECT_SOURCE_DIR}" STREQUAL "${PROJECT_BINARY_DIR}") + message(STATUS "") + message(STATUS "rocprofiler does not support in-source builds.") + message(STATUS "Delete CMakeCache.txt and CMakeFiles in ${PROJECT_SOURCE_DIR}") + message(STATUS "and run cmake with `-B `") + message(STATUS "") + message(FATAL_ERROR "In-source build detected.") +endif() + +option(ROCPROFILER_BUILD_TESTS "Enable building the tests" OFF) +option(ROCPROFILER_BUILD_SAMPLES "Enable building the code samples" OFF) + +# CLI and FILE plugins are always built +foreach(_PLUGIN "ATT" "CTF" "PERFETTO") + option(ROCPROFILER_BUILD_PLUGIN_${_PLUGIN} "Enable building the ${_PLUGIN} plugin" ON) +endforeach() + +option(ROCPROFILER_DEBUG_TRACE "Enable debug tracing" OFF) +mark_as_advanced(ROCPROFILER_DEBUG_TRACE) + +option(ROCPROFILER_LD_AQLPROFILE "Enable direct loading of AQL-profile HSA extension" OFF) +mark_as_advanced(ROCPROFILER_LD_AQLPROFILE) + +option(ROCPROFILER_BUILD_CI "Enable continuous integration additions" OFF) +mark_as_advanced(ROCPROFILER_BUILD_CI) + +option(ROCPROFILER_ENABLE_CLANG_TIDY "Enable clang-tidy checks" OFF) +mark_as_advanced(ROCPROFILER_ENABLE_CLANG_TIDY) + +set(ROCPROFILER_BUILD_TYPES "Release" "RelWithDebInfo" "Debug" "MinSizeRel" "Coverage") + +# export compile commands in the project +set(CMAKE_EXPORT_COMPILE_COMMANDS ON) + +if(NOT CMAKE_BUILD_TYPE) + set(CMAKE_BUILD_TYPE + "Release" + CACHE STRING "Build type" FORCE) +endif() + +if(NOT CMAKE_BUILD_TYPE IN_LIST ROCPROFILER_BUILD_TYPES) + message( + FATAL_ERROR + "Unsupported build type '${CMAKE_BUILD_TYPE}'. Options: ${ROCPROFILER_BUILD_TYPES}" + ) +endif() + +if(ROCPROFILER_BUILD_CI) + foreach(_BUILD_TYPE ${ROCPROFILER_BUILD_TYPES}) + string(TOUPPER "${_BUILD_TYPE}" _BUILD_TYPE) + + # remove NDEBUG preprocessor def so that asserts are triggered + string(REGEX REPLACE ".DNDEBUG" "" CMAKE_C_FLAGS_${_BUILD_TYPE} + "${CMAKE_C_FLAGS_${_BUILD_TYPE}}") + string(REGEX REPLACE ".DNDEBUG" "" CMAKE_CXX_FLAGS_${_BUILD_TYPE} + "${CMAKE_CXX_FLAGS_${_BUILD_TYPE}}") + endforeach() +endif() + +if(CMAKE_PROJECT_NAME STREQUAL PROJECT_NAME) + set_property(CACHE CMAKE_BUILD_TYPE PROPERTY STRINGS "${ROCPROFILER_BUILD_TYPES}") +endif() + +set(ROCPROFILER_MEMCHECK + "" + CACHE STRING "Memory checker type") +mark_as_advanced(ROCPROFILER_MEMCHECK) + +# ASAN is defined by testing team on Jenkins +if(ASAN) + set(ROCPROFILER_MEMCHECK + "AddressSanitizer" + CACHE STRING "Memory checker type (forced by ASAN defined)" FORCE) +endif() + +set(ROCPROFILER_MEMCHECK_TYPES "ThreadSanitizer" "AddressSanitizer" "LeakSanitizer" + "MemorySanitizer" "UndefinedBehaviorSanitizer") + +if(ROCPROFILER_MEMCHECK AND NOT ROCPROFILER_MEMCHECK IN_LIST ROCPROFILER_MEMCHECK_TYPES) + message( + FATAL_ERROR + "Unsupported memcheck type '${ROCPROFILER_MEMCHECK}'. Options: ${ROCPROFILER_MEMCHECK_TYPES}" + ) +endif() + +set_property(CACHE ROCPROFILER_MEMCHECK PROPERTY STRINGS "${ROCPROFILER_MEMCHECK_TYPES}") + +add_library(rocprofiler-memcheck INTERFACE) +add_library(rocprofiler::memcheck ALIAS rocprofiler-memcheck) + +function(rocprofiler_add_memcheck_flags _TYPE) + target_compile_options( + rocprofiler-memcheck INTERFACE $) + target_link_options(rocprofiler-memcheck INTERFACE + $) +endfunction() + +function(rocprofiler_set_memcheck_env _TYPE _LIB_BASE) + set(_LIBS ${_LIB_BASE}) + foreach(_N 6 5 4 3 2 1 0) + list( + APPEND _LIBS + ${CMAKE_SHARED_LIBRARY_PREFIX}${_LIB_BASE}${CMAKE_SHARED_LIBRARY_SUFFIX}.${_N} + ) + endforeach() + foreach(_LIB ${_LIBS}) + if(NOT ${_TYPE}_LIBRARY) + find_library(${_TYPE}_LIBRARY NAMES ${_LIB} ${ARGN}) + endif() + endforeach() + + target_link_libraries(rocprofiler-memcheck INTERFACE ${_LIB_BASE}) + + if(${_TYPE}_LIBRARY) + set(ROCPROFILER_MEMCHECK_PRELOAD_ENV + "LD_PRELOAD=${${_TYPE}_LIBRARY};LD_LIBRARY_PATH=${PROJECT_BINARY_DIR}/lib:$ENV{LD_LIBRARY_PATH}" + CACHE INTERNAL "LD_PRELOAD env variable for tests" FORCE) + endif() +endfunction() + +# always unset so that it doesn't preload if memcheck disabled +unset(ROCPROFILER_MEMCHECK_PRELOAD_ENV CACHE) + +if(ROCPROFILER_MEMCHECK STREQUAL "AddressSanitizer") + rocprofiler_add_memcheck_flags("address") + rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "asan") +elseif(ROCPROFILER_MEMCHECK STREQUAL "LeakSanitizer") + rocprofiler_add_memcheck_flags("leak") + rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "lsan") +elseif(ROCPROFILER_MEMCHECK STREQUAL "MemorySanitizer") + rocprofiler_add_memcheck_flags("memory") +elseif(ROCPROFILER_MEMCHECK STREQUAL "ThreadSanitizer") + rocprofiler_add_memcheck_flags("thread") + rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "tsan") +elseif(ROCPROFILER_MEMCHECK STREQUAL "UndefinedBehaviorSanitizer") + rocprofiler_add_memcheck_flags("undefined") + rocprofiler_set_memcheck_env("${ROCPROFILER_MEMCHECK}" "ubsan") +endif() diff --git a/projects/rocprofiler/cmake_modules/rocprofiler_utils.cmake b/projects/rocprofiler/cmake_modules/rocprofiler_utils.cmake new file mode 100644 index 0000000000..ad4134eb2d --- /dev/null +++ b/projects/rocprofiler/cmake_modules/rocprofiler_utils.cmake @@ -0,0 +1,263 @@ +################################################################################ +# Copyright (c) 2018 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. +################################################################################ + +# Parses the VERSION_STRING variable and places the first, second and third number values +# in the major, minor and patch variables. +function(rocprofiler_parse_version VERSION_STRING) + + string(FIND ${VERSION_STRING} "-" STRING_INDEX) + + if(${STRING_INDEX} GREATER -1) + math(EXPR STRING_INDEX "${STRING_INDEX} + 1") + string(SUBSTRING ${VERSION_STRING} ${STRING_INDEX} -1 VERSION_BUILD) + endif() + + string(REGEX MATCHALL "[0123456789]+" VERSIONS ${VERSION_STRING}) + list(LENGTH VERSIONS VERSION_COUNT) + + if(${VERSION_COUNT} GREATER 0) + list(GET VERSIONS 0 MAJOR) + set(VERSION_MAJOR + ${MAJOR} + PARENT_SCOPE) + set(TEMP_VERSION_STRING "${MAJOR}") + endif() + + if(${VERSION_COUNT} GREATER 1) + list(GET VERSIONS 1 MINOR) + set(VERSION_MINOR + ${MINOR} + PARENT_SCOPE) + set(TEMP_VERSION_STRING "${TEMP_VERSION_STRING}.${MINOR}") + endif() + + if(${VERSION_COUNT} GREATER 2) + list(GET VERSIONS 2 PATCH) + set(VERSION_PATCH + ${PATCH} + PARENT_SCOPE) + set(TEMP_VERSION_STRING "${TEMP_VERSION_STRING}.${PATCH}") + endif() + + if(DEFINED VERSION_BUILD) + set(VERSION_BUILD + "${VERSION_BUILD}" + PARENT_SCOPE) + endif() + + set(VERSION_STRING + "${TEMP_VERSION_STRING}" + PARENT_SCOPE) + +endfunction() + +# Gets the current version of the repository using versioning tags and git describe. +# Passes back a packaging version string and a library version string. +function(rocprofiler_get_version DEFAULT_VERSION_STRING) + + rocprofiler_parse_version(${DEFAULT_VERSION_STRING}) + + find_program(GIT NAMES git) + + if(GIT) + + execute_process( + COMMAND "git describe --dirty --long --match [0-9]* 2>/dev/null" + WORKING_DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} + OUTPUT_VARIABLE GIT_TAG_STRING + OUTPUT_STRIP_TRAILING_WHITESPACE + RESULT_VARIABLE RESULT) + + if(${RESULT} EQUAL 0) + + rocprofiler_parse_version(${GIT_TAG_STRING}) + + endif() + + endif() + + set(VERSION_STRING + "${VERSION_STRING}" + PARENT_SCOPE) + set(VERSION_MAJOR + "${VERSION_MAJOR}" + PARENT_SCOPE) + set(VERSION_MINOR + "${VERSION_MINOR}" + PARENT_SCOPE) + set(VERSION_PATCH + "${VERSION_PATCH}" + PARENT_SCOPE) + set(VERSION_BUILD + "${VERSION_BUILD}" + PARENT_SCOPE) + +endfunction() + +# ----------------------------------------------------------------------------------------# +# macro rocprofiler_checkout_git_submodule() +# +# Run "git submodule update" if a file in a submodule does not exist +# +# ARGS: RECURSIVE (option) -- add "--recursive" flag RELATIVE_PATH (one value) -- +# typically the relative path to submodule from PROJECT_SOURCE_DIR WORKING_DIRECTORY (one +# value) -- (default: PROJECT_SOURCE_DIR) TEST_FILE (one value) -- file to check for +# (default: CMakeLists.txt) ADDITIONAL_CMDS (many value) -- any addition commands to pass +# +function(ROCPROFILER_CHECKOUT_GIT_SUBMODULE) + # parse args + cmake_parse_arguments( + CHECKOUT "RECURSIVE" + "RELATIVE_PATH;WORKING_DIRECTORY;TEST_FILE;REPO_URL;REPO_BRANCH" + "ADDITIONAL_CMDS" ${ARGN}) + + if(NOT CHECKOUT_WORKING_DIRECTORY) + set(CHECKOUT_WORKING_DIRECTORY ${PROJECT_SOURCE_DIR}) + endif() + + if(NOT CHECKOUT_TEST_FILE) + set(CHECKOUT_TEST_FILE "CMakeLists.txt") + endif() + + find_package(Git) + set(_DIR "${CHECKOUT_WORKING_DIRECTORY}/${CHECKOUT_RELATIVE_PATH}") + # ensure the (possibly empty) directory exists + if(NOT EXISTS "${_DIR}") + if(NOT CHECKOUT_REPO_URL) + message(FATAL_ERROR "submodule directory does not exist") + endif() + endif() + + # if this file exists --> project has been checked out if not exists --> not been + # checked out + set(_TEST_FILE "${_DIR}/${CHECKOUT_TEST_FILE}") + # assuming a .gitmodules file exists + set(_SUBMODULE "${PROJECT_SOURCE_DIR}/.gitmodules") + + set(_TEST_FILE_EXISTS OFF) + if(EXISTS "${_TEST_FILE}" AND NOT IS_DIRECTORY "${_TEST_FILE}") + set(_TEST_FILE_EXISTS ON) + endif() + + if(_TEST_FILE_EXISTS) + return() + endif() + + find_package(Git REQUIRED) + + set(_SUBMODULE_EXISTS OFF) + if(EXISTS "${_SUBMODULE}" AND NOT IS_DIRECTORY "${_SUBMODULE}") + set(_SUBMODULE_EXISTS ON) + endif() + + set(_HAS_REPO_URL OFF) + if(NOT "${CHECKOUT_REPO_URL}" STREQUAL "") + set(_HAS_REPO_URL ON) + endif() + + # if the module has not been checked out + if(NOT _TEST_FILE_EXISTS AND _SUBMODULE_EXISTS) + # perform the checkout + execute_process( + COMMAND ${GIT_EXECUTABLE} submodule update --init ${_RECURSE} + ${CHECKOUT_ADDITIONAL_CMDS} ${CHECKOUT_RELATIVE_PATH} + WORKING_DIRECTORY ${CHECKOUT_WORKING_DIRECTORY} + RESULT_VARIABLE RET) + + # check the return code + if(RET GREATER 0) + set(_CMD "${GIT_EXECUTABLE} submodule update --init ${_RECURSE} + ${CHECKOUT_ADDITIONAL_CMDS} ${CHECKOUT_RELATIVE_PATH}") + message(STATUS "function(rocprofiler_checkout_git_submodule) failed.") + message(FATAL_ERROR "Command: \"${_CMD}\"") + else() + set(_TEST_FILE_EXISTS ON) + endif() + endif() + + if(NOT _TEST_FILE_EXISTS AND _HAS_REPO_URL) + message( + STATUS + "Cloning '${CHECKOUT_REPO_URL}' into ${CHECKOUT_WORKING_DIRECTORY}/${CHECKOUT_RELATIVE_PATH}..." + ) + + # remove the existing directory + if(EXISTS "${_DIR}") + execute_process(COMMAND ${CMAKE_COMMAND} -E remove_directory ${_DIR}) + endif() + + # perform the checkout + execute_process( + COMMAND ${GIT_EXECUTABLE} clone ${CHECKOUT_ADDITIONAL_CMDS} + ${CHECKOUT_REPO_URL} ${CHECKOUT_RELATIVE_PATH} + WORKING_DIRECTORY ${CHECKOUT_WORKING_DIRECTORY} + RESULT_VARIABLE RET_CLONE) + + if(NOT RET_CLONE EQUAL 0) + message( + SEND_ERROR + "Failed to clone ${CHECKOUT_REPO_URL} into ${CHECKOUT_WORKING_DIRECTORY}/${CHECKOUT_RELATIVE_PATH}" + ) + return() + endif() + + if(CHECKOUT_REPO_BRANCH) + execute_process( + COMMAND ${GIT_EXECUTABLE} checkout ${CHECKOUT_REPO_BRANCH} + WORKING_DIRECTORY ${CHECKOUT_WORKING_DIRECTORY}/${CHECKOUT_RELATIVE_PATH} + RESULT_VARIABLE RET_BRANCH) + + if(NOT RET_BRANCH EQUAL 0) + message( + SEND_ERROR + "Failed to checkout '${CHECKOUT_REPO_BRANCH}' for ${CHECKOUT_REPO_URL} in ${CHECKOUT_WORKING_DIRECTORY}/${CHECKOUT_RELATIVE_PATH}" + ) + return() + endif() + endif() + + # perform the submodule update + if(CHECKOUT_RECURSIVE + AND EXISTS "${_DIR}" + AND IS_DIRECTORY "${_DIR}") + execute_process( + COMMAND ${GIT_EXECUTABLE} submodule update --init ${_RECURSE} + WORKING_DIRECTORY ${_DIR} + RESULT_VARIABLE RET_RECURSIVE) + if(NOT RET_RECURSIVE EQUAL 0) + message( + SEND_ERROR + "Failed to update submodules for ${CHECKOUT_REPO_URL} in ${CHECKOUT_WORKING_DIRECTORY}/${CHECKOUT_RELATIVE_PATH}" + ) + return() + endif() + endif() + + set(_TEST_FILE_EXISTS ON) + endif() + + if(NOT EXISTS "${_TEST_FILE}" OR NOT _TEST_FILE_EXISTS) + message( + FATAL_ERROR + "Error checking out submodule: '${CHECKOUT_RELATIVE_PATH}' to '${_DIR}'") + endif() +endfunction() diff --git a/projects/rocprofiler/doc/Doxyfile_API.in b/projects/rocprofiler/doc/Doxyfile_API.in new file mode 100644 index 0000000000..c94a585a3a --- /dev/null +++ b/projects/rocprofiler/doc/Doxyfile_API.in @@ -0,0 +1,2447 @@ +## Copyright (c) 2018-2022 Advanced Micro Devices, Inc. +## +## 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. + +# Doxyfile 1.8.11 + +# This file describes the settings to be used by the documentation system +# doxygen (www.doxygen.org) for a project. +# +# All text after a double hash (##) is considered a comment and is placed in +# front of the TAG it is preceding. +# +# All text after a single hash (#) is considered a comment and will be ignored. +# The format is: +# TAG = value [value, ...] +# For lists, items can also be appended using: +# TAG += value [value, ...] +# Values that contain spaces should be placed between quotes (\" \"). + +#--------------------------------------------------------------------------- +# Project related configuration options +#--------------------------------------------------------------------------- + +# This tag specifies the encoding used for all characters in the config file +# that follow. The default is UTF-8 which is also the encoding used for all text +# before the first occurrence of this tag. Doxygen uses libiconv (or the iconv +# built into libc) for the transcoding. See http://www.gnu.org/software/libiconv +# for the list of possible encodings. +# The default value is: UTF-8. + +DOXYFILE_ENCODING = UTF-8 + +# The PROJECT_NAME tag is a single word (or a sequence of words surrounded by +# double-quotes, unless you are using Doxywizard) that should identify the +# project for which the documentation is generated. This name is used in the +# title of most generated pages and in a few other places. +# The default value is: My Project. + +PROJECT_NAME = "@PROJECT_NAME@" + +# The PROJECT_NUMBER tag can be used to enter a project or revision number. This +# could be handy for archiving the generated documentation or if some version +# control system is used. + +PROJECT_NUMBER = @PROJECT_VERSION_MAJOR@.@PROJECT_VERSION_MINOR@.@PROJECT_VERSION_PATCH@ + +# Using the PROJECT_BRIEF tag one can provide an optional one line description +# for a project that appears at the top of each page and should give viewer a +# quick idea about the purpose of the project. Keep the description short. + +PROJECT_BRIEF = + +# With the PROJECT_LOGO tag one can specify a logo or an icon that is included +# in the documentation. The maximum height of the logo should not exceed 55 +# pixels and the maximum width should not exceed 200 pixels. Doxygen will copy +# the logo to the output directory. + +PROJECT_LOGO = + +# The OUTPUT_DIRECTORY tag is used to specify the (relative or absolute) path +# into which the generated documentation will be written. If a relative path is +# entered, it will be relative to the location where doxygen was started. If +# left blank the current directory will be used. + +OUTPUT_DIRECTORY = @CMAKE_CURRENT_BINARY_DIR@/doc/ + +# If the CREATE_SUBDIRS tag is set to YES then doxygen will create 4096 sub- +# directories (in 2 levels) under the output directory of each output format and +# will distribute the generated files over these directories. Enabling this +# option can be useful when feeding doxygen a huge amount of source files, where +# putting all generated files in the same directory would otherwise causes +# performance problems for the file system. +# The default value is: NO. + +CREATE_SUBDIRS = NO + +# If the ALLOW_UNICODE_NAMES tag is set to YES, doxygen will allow non-ASCII +# characters to appear in the names of generated files. If set to NO, non-ASCII +# characters will be escaped, for example _xE3_x81_x84 will be used for Unicode +# U+3044. +# The default value is: NO. + +ALLOW_UNICODE_NAMES = NO + +# The OUTPUT_LANGUAGE tag is used to specify the language in which all +# documentation generated by doxygen is written. Doxygen will use this +# information to generate all constant output in the proper language. +# Possible values are: Afrikaans, Arabic, Armenian, Brazilian, Catalan, Chinese, +# Chinese-Traditional, Croatian, Czech, Danish, Dutch, English (United States), +# Esperanto, Farsi (Persian), Finnish, French, German, Greek, Hungarian, +# Indonesian, Italian, Japanese, Japanese-en (Japanese with English messages), +# Korean, Korean-en (Korean with English messages), Latvian, Lithuanian, +# Macedonian, Norwegian, Persian (Farsi), Polish, Portuguese, Romanian, Russian, +# Serbian, Serbian-Cyrillic, Slovak, Slovene, Spanish, Swedish, Turkish, +# Ukrainian and Vietnamese. +# The default value is: English. + +OUTPUT_LANGUAGE = English + +# If the BRIEF_MEMBER_DESC tag is set to YES, doxygen will include brief member +# descriptions after the members that are listed in the file and class +# documentation (similar to Javadoc). Set to NO to disable this. +# The default value is: YES. + +BRIEF_MEMBER_DESC = YES + +# If the REPEAT_BRIEF tag is set to YES, doxygen will prepend the brief +# description of a member or function before the detailed description +# +# Note: If both HIDE_UNDOC_MEMBERS and BRIEF_MEMBER_DESC are set to NO, the +# brief descriptions will be completely suppressed. +# The default value is: YES. + +REPEAT_BRIEF = YES + +# This tag implements a quasi-intelligent brief description abbreviator that is +# used to form the text in various listings. Each string in this list, if found +# as the leading text of the brief description, will be stripped from the text +# and the result, after processing the whole list, is used as the annotated +# text. Otherwise, the brief description is used as-is. If left blank, the +# following values are used ($name is automatically replaced with the name of +# the entity):The $name class, The $name widget, The $name file, is, provides, +# specifies, contains, represents, a, an and the. + +ABBREVIATE_BRIEF = + +# If the ALWAYS_DETAILED_SEC and REPEAT_BRIEF tags are both set to YES then +# doxygen will generate a detailed section even if there is only a brief +# description. +# The default value is: NO. + +ALWAYS_DETAILED_SEC = NO + +# If the INLINE_INHERITED_MEMB tag is set to YES, doxygen will show all +# inherited members of a class in the documentation of that class as if those +# members were ordinary class members. Constructors, destructors and assignment +# operators of the base classes will not be shown. +# The default value is: NO. + +INLINE_INHERITED_MEMB = NO + +# If the FULL_PATH_NAMES tag is set to YES, doxygen will prepend the full path +# before files name in the file list and in the header files. If set to NO the +# shortest path that makes the file name unique will be used +# The default value is: YES. + +FULL_PATH_NAMES = YES + +# The STRIP_FROM_PATH tag can be used to strip a user-defined part of the path. +# Stripping is only done if one of the specified strings matches the left-hand +# part of the path. The tag can be used to show relative paths in the file list. +# If left blank the directory from which doxygen is run is used as the path to +# strip. +# +# Note that you can specify absolute paths here, but also relative paths, which +# will be relative from the directory where doxygen is started. +# This tag requires that the tag FULL_PATH_NAMES is set to YES. + +STRIP_FROM_PATH = + +# The STRIP_FROM_INC_PATH tag can be used to strip a user-defined part of the +# path mentioned in the documentation of a class, which tells the reader which +# header file to include in order to use a class. If left blank only the name of +# the header file containing the class definition is used. Otherwise one should +# specify the list of include paths that are normally passed to the compiler +# using the -I flag. + +STRIP_FROM_INC_PATH = + +# If the SHORT_NAMES tag is set to YES, doxygen will generate much shorter (but +# less readable) file names. This can be useful is your file systems doesn't +# support long names like on DOS, Mac, or CD-ROM. +# The default value is: NO. + +SHORT_NAMES = NO + +# If the JAVADOC_AUTOBRIEF tag is set to YES then doxygen will interpret the +# first line (until the first dot) of a Javadoc-style comment as the brief +# description. If set to NO, the Javadoc-style will behave just like regular Qt- +# style comments (thus requiring an explicit @brief command for a brief +# description.) +# The default value is: NO. + +JAVADOC_AUTOBRIEF = YES + +# If the QT_AUTOBRIEF tag is set to YES then doxygen will interpret the first +# line (until the first dot) of a Qt-style comment as the brief description. If +# set to NO, the Qt-style will behave just like regular Qt-style comments (thus +# requiring an explicit \brief command for a brief description.) +# The default value is: NO. + +QT_AUTOBRIEF = NO + +# The MULTILINE_CPP_IS_BRIEF tag can be set to YES to make doxygen treat a +# multi-line C++ special comment block (i.e. a block of //! or /// comments) as +# a brief description. This used to be the default behavior. The new default is +# to treat a multi-line C++ comment block as a detailed description. Set this +# tag to YES if you prefer the old behavior instead. +# +# Note that setting this tag to YES also means that rational rose comments are +# not recognized any more. +# The default value is: NO. + +MULTILINE_CPP_IS_BRIEF = NO + +# If the INHERIT_DOCS tag is set to YES then an undocumented member inherits the +# documentation from any documented member that it re-implements. +# The default value is: YES. + +INHERIT_DOCS = YES + +# If the SEPARATE_MEMBER_PAGES tag is set to YES then doxygen will produce a new +# page for each member. If set to NO, the documentation of a member will be part +# of the file/class/namespace that contains it. +# The default value is: NO. + +SEPARATE_MEMBER_PAGES = NO + +# The TAB_SIZE tag can be used to set the number of spaces in a tab. Doxygen +# uses this value to replace tabs by spaces in code fragments. +# Minimum value: 1, maximum value: 16, default value: 4. + +TAB_SIZE = 4 + +# This tag can be used to specify a number of aliases that act as commands in +# the documentation. An alias has the form: +# name=value +# For example adding +# "sideeffect=@par Side Effects:\n" +# will allow you to put the command \sideeffect (or @sideeffect) in the +# documentation, which will result in a user-defined paragraph with heading +# "Side Effects:". You can put \n's in the value part of an alias to insert +# newlines. + +ALIASES = + +# This tag can be used to specify a number of word-keyword mappings (TCL only). +# A mapping has the form "name=value". For example adding "class=itcl::class" +# will allow you to use the command class in the itcl::class meaning. + +TCL_SUBST = + +# Set the OPTIMIZE_OUTPUT_FOR_C tag to YES if your project consists of C sources +# only. Doxygen will then generate output that is more tailored for C. For +# instance, some of the names that are used will be different. The list of all +# members will be omitted, etc. +# The default value is: NO. + +OPTIMIZE_OUTPUT_FOR_C = YES + +# Set the OPTIMIZE_OUTPUT_JAVA tag to YES if your project consists of Java or +# Python sources only. Doxygen will then generate output that is more tailored +# for that language. For instance, namespaces will be presented as packages, +# qualified scopes will look different, etc. +# The default value is: NO. + +OPTIMIZE_OUTPUT_JAVA = NO + +# Set the OPTIMIZE_FOR_FORTRAN tag to YES if your project consists of Fortran +# sources. Doxygen will then generate output that is tailored for Fortran. +# The default value is: NO. + +OPTIMIZE_FOR_FORTRAN = NO + +# Set the OPTIMIZE_OUTPUT_VHDL tag to YES if your project consists of VHDL +# sources. Doxygen will then generate output that is tailored for VHDL. +# The default value is: NO. + +OPTIMIZE_OUTPUT_VHDL = NO + +# Doxygen selects the parser to use depending on the extension of the files it +# parses. With this tag you can assign which parser to use for a given +# extension. Doxygen has a built-in mapping, but you can override or extend it +# using this tag. The format is ext=language, where ext is a file extension, and +# language is one of the parsers supported by doxygen: IDL, Java, Javascript, +# C#, C, C++, D, PHP, Objective-C, Python, Fortran (fixed format Fortran: +# FortranFixed, free formatted Fortran: FortranFree, unknown formatted Fortran: +# Fortran. In the later case the parser tries to guess whether the code is fixed +# or free formatted code, this is the default for Fortran type files), VHDL. For +# instance to make doxygen treat .inc files as Fortran files (default is PHP), +# and .f files as C (default is Fortran), use: inc=Fortran f=C. +# +# Note: For files without extension you can use no_extension as a placeholder. +# +# Note that for custom extensions you also need to set FILE_PATTERNS otherwise +# the files are not read by doxygen. + +EXTENSION_MAPPING = + +# If the MARKDOWN_SUPPORT tag is enabled then doxygen pre-processes all comments +# according to the Markdown format, which allows for more readable +# documentation. See http://daringfireball.net/projects/markdown/ for details. +# The output of markdown processing is further processed by doxygen, so you can +# mix doxygen, HTML, and XML commands with Markdown formatting. Disable only in +# case of backward compatibilities issues. +# The default value is: YES. + +MARKDOWN_SUPPORT = YES + +# When enabled doxygen tries to link words that correspond to documented +# classes, or namespaces to their corresponding documentation. Such a link can +# be prevented in individual cases by putting a % sign in front of the word or +# globally by setting AUTOLINK_SUPPORT to NO. +# The default value is: YES. + +AUTOLINK_SUPPORT = YES + +# If you use STL classes (i.e. std::string, std::vector, etc.) but do not want +# to include (a tag file for) the STL sources as input, then you should set this +# tag to YES in order to let doxygen match functions declarations and +# definitions whose arguments contain STL classes (e.g. func(std::string); +# versus func(std::string) {}). This also make the inheritance and collaboration +# diagrams that involve STL classes more complete and accurate. +# The default value is: NO. + +BUILTIN_STL_SUPPORT = NO + +# If you use Microsoft's C++/CLI language, you should set this option to YES to +# enable parsing support. +# The default value is: NO. + +CPP_CLI_SUPPORT = NO + +# Set the SIP_SUPPORT tag to YES if your project consists of sip (see: +# http://www.riverbankcomputing.co.uk/software/sip/intro) sources only. Doxygen +# will parse them like normal C++ but will assume all classes use public instead +# of private inheritance when no explicit protection keyword is present. +# The default value is: NO. + +SIP_SUPPORT = NO + +# For Microsoft's IDL there are propget and propput attributes to indicate +# getter and setter methods for a property. Setting this option to YES will make +# doxygen to replace the get and set methods by a property in the documentation. +# This will only work if the methods are indeed getting or setting a simple +# type. If this is not the case, or you want to show the methods anyway, you +# should set this option to NO. +# The default value is: YES. + +IDL_PROPERTY_SUPPORT = YES + +# If member grouping is used in the documentation and the DISTRIBUTE_GROUP_DOC +# tag is set to YES then doxygen will reuse the documentation of the first +# member in the group (if any) for the other members of the group. By default +# all members of a group must be documented explicitly. +# The default value is: NO. + +DISTRIBUTE_GROUP_DOC = NO + +# If one adds a struct or class to a group and this option is enabled, then also +# any nested class or struct is added to the same group. By default this option +# is disabled and one has to add nested compounds explicitly via \ingroup. +# The default value is: NO. + +GROUP_NESTED_COMPOUNDS = NO + +# Set the SUBGROUPING tag to YES to allow class member groups of the same type +# (for instance a group of public functions) to be put as a subgroup of that +# type (e.g. under the Public Functions section). Set it to NO to prevent +# subgrouping. Alternatively, this can be done per class using the +# \nosubgrouping command. +# The default value is: YES. + +SUBGROUPING = YES + +# When the INLINE_GROUPED_CLASSES tag is set to YES, classes, structs and unions +# are shown inside the group in which they are included (e.g. using \ingroup) +# instead of on a separate page (for HTML and Man pages) or section (for LaTeX +# and RTF). +# +# Note that this feature does not work in combination with +# SEPARATE_MEMBER_PAGES. +# The default value is: NO. + +INLINE_GROUPED_CLASSES = NO + +# When the INLINE_SIMPLE_STRUCTS tag is set to YES, structs, classes, and unions +# with only public data fields or simple typedef fields will be shown inline in +# the documentation of the scope in which they are defined (i.e. file, +# namespace, or group documentation), provided this scope is documented. If set +# to NO, structs, classes, and unions are shown on a separate page (for HTML and +# Man pages) or section (for LaTeX and RTF). +# The default value is: NO. + +INLINE_SIMPLE_STRUCTS = NO + +# When TYPEDEF_HIDES_STRUCT tag is enabled, a typedef of a struct, union, or +# enum is documented as struct, union, or enum with the name of the typedef. So +# typedef struct TypeS {} TypeT, will appear in the documentation as a struct +# with name TypeT. When disabled the typedef will appear as a member of a file, +# namespace, or class. And the struct will be named TypeS. This can typically be +# useful for C code in case the coding convention dictates that all compound +# types are typedef'ed and only the typedef is referenced, never the tag name. +# The default value is: NO. + +TYPEDEF_HIDES_STRUCT = NO + +# The size of the symbol lookup cache can be set using LOOKUP_CACHE_SIZE. This +# cache is used to resolve symbols given their name and scope. Since this can be +# an expensive process and often the same symbol appears multiple times in the +# code, doxygen keeps a cache of pre-resolved symbols. If the cache is too small +# doxygen will become slower. If the cache is too large, memory is wasted. The +# cache size is given by this formula: 2^(16+LOOKUP_CACHE_SIZE). The valid range +# is 0..9, the default is 0, corresponding to a cache size of 2^16=65536 +# symbols. At the end of a run doxygen will report the cache usage and suggest +# the optimal cache size from a speed point of view. +# Minimum value: 0, maximum value: 9, default value: 0. + +LOOKUP_CACHE_SIZE = 0 + +#--------------------------------------------------------------------------- +# Build related configuration options +#--------------------------------------------------------------------------- + +# If the EXTRACT_ALL tag is set to YES, doxygen will assume all entities in +# documentation are documented, even if no documentation was available. Private +# class members and static file members will be hidden unless the +# EXTRACT_PRIVATE respectively EXTRACT_STATIC tags are set to YES. +# Note: This will also disable the warnings about undocumented members that are +# normally produced when WARNINGS is set to YES. +# The default value is: NO. + +EXTRACT_ALL = YES + +# If the EXTRACT_PRIVATE tag is set to YES, all private members of a class will +# be included in the documentation. +# The default value is: NO. + +EXTRACT_PRIVATE = NO + +# If the EXTRACT_PACKAGE tag is set to YES, all members with package or internal +# scope will be included in the documentation. +# The default value is: NO. + +EXTRACT_PACKAGE = NO + +# If the EXTRACT_STATIC tag is set to YES, all static members of a file will be +# included in the documentation. +# The default value is: NO. + +EXTRACT_STATIC = NO + +# If the EXTRACT_LOCAL_CLASSES tag is set to YES, classes (and structs) defined +# locally in source files will be included in the documentation. If set to NO, +# only classes defined in header files are included. Does not have any effect +# for Java sources. +# The default value is: YES. + +EXTRACT_LOCAL_CLASSES = YES + +# This flag is only useful for Objective-C code. If set to YES, local methods, +# which are defined in the implementation section but not in the interface are +# included in the documentation. If set to NO, only methods in the interface are +# included. +# The default value is: NO. + +EXTRACT_LOCAL_METHODS = NO + +# If this flag is set to YES, the members of anonymous namespaces will be +# extracted and appear in the documentation as a namespace called +# 'anonymous_namespace{file}', where file will be replaced with the base name of +# the file that contains the anonymous namespace. By default anonymous namespace +# are hidden. +# The default value is: NO. + +EXTRACT_ANON_NSPACES = NO + +# If the HIDE_UNDOC_MEMBERS tag is set to YES, doxygen will hide all +# undocumented members inside documented classes or files. If set to NO these +# members will be included in the various overviews, but no documentation +# section is generated. This option has no effect if EXTRACT_ALL is enabled. +# The default value is: NO. + +HIDE_UNDOC_MEMBERS = NO + +# If the HIDE_UNDOC_CLASSES tag is set to YES, doxygen will hide all +# undocumented classes that are normally visible in the class hierarchy. If set +# to NO, these classes will be included in the various overviews. This option +# has no effect if EXTRACT_ALL is enabled. +# The default value is: NO. + +HIDE_UNDOC_CLASSES = NO + +# If the HIDE_FRIEND_COMPOUNDS tag is set to YES, doxygen will hide all friend +# (class|struct|union) declarations. If set to NO, these declarations will be +# included in the documentation. +# The default value is: NO. + +HIDE_FRIEND_COMPOUNDS = NO + +# If the HIDE_IN_BODY_DOCS tag is set to YES, doxygen will hide any +# documentation blocks found inside the body of a function. If set to NO, these +# blocks will be appended to the function's detailed documentation block. +# The default value is: NO. + +HIDE_IN_BODY_DOCS = NO + +# The INTERNAL_DOCS tag determines if documentation that is typed after a +# \internal command is included. If the tag is set to NO then the documentation +# will be excluded. Set it to YES to include the internal documentation. +# The default value is: NO. + +INTERNAL_DOCS = NO + +# If the CASE_SENSE_NAMES tag is set to NO then doxygen will only generate file +# names in lower-case letters. If set to YES, upper-case letters are also +# allowed. This is useful if you have classes or files whose names only differ +# in case and if your file system supports case sensitive file names. Windows +# and Mac users are advised to set this option to NO. +# The default value is: system dependent. + +CASE_SENSE_NAMES = YES + +# If the HIDE_SCOPE_NAMES tag is set to NO then doxygen will show members with +# their full class and namespace scopes in the documentation. If set to YES, the +# scope will be hidden. +# The default value is: NO. + +HIDE_SCOPE_NAMES = NO + +# If the HIDE_COMPOUND_REFERENCE tag is set to NO (default) then doxygen will +# append additional text to a page's title, such as Class Reference. If set to +# YES the compound reference will be hidden. +# The default value is: NO. + +HIDE_COMPOUND_REFERENCE= NO + +# If the SHOW_INCLUDE_FILES tag is set to YES then doxygen will put a list of +# the files that are included by a file in the documentation of that file. +# The default value is: YES. + +SHOW_INCLUDE_FILES = YES + +# If the SHOW_GROUPED_MEMB_INC tag is set to YES then Doxygen will add for each +# grouped member an include statement to the documentation, telling the reader +# which file to include in order to use the member. +# The default value is: NO. + +SHOW_GROUPED_MEMB_INC = NO + +# If the FORCE_LOCAL_INCLUDES tag is set to YES then doxygen will list include +# files with double quotes in the documentation rather than with sharp brackets. +# The default value is: NO. + +FORCE_LOCAL_INCLUDES = NO + +# If the INLINE_INFO tag is set to YES then a tag [inline] is inserted in the +# documentation for inline members. +# The default value is: YES. + +INLINE_INFO = YES + +# If the SORT_MEMBER_DOCS tag is set to YES then doxygen will sort the +# (detailed) documentation of file and class members alphabetically by member +# name. If set to NO, the members will appear in declaration order. +# The default value is: YES. + +SORT_MEMBER_DOCS = YES + +# If the SORT_BRIEF_DOCS tag is set to YES then doxygen will sort the brief +# descriptions of file, namespace and class members alphabetically by member +# name. If set to NO, the members will appear in declaration order. Note that +# this will also influence the order of the classes in the class list. +# The default value is: NO. + +SORT_BRIEF_DOCS = NO + +# If the SORT_MEMBERS_CTORS_1ST tag is set to YES then doxygen will sort the +# (brief and detailed) documentation of class members so that constructors and +# destructors are listed first. If set to NO the constructors will appear in the +# respective orders defined by SORT_BRIEF_DOCS and SORT_MEMBER_DOCS. +# Note: If SORT_BRIEF_DOCS is set to NO this option is ignored for sorting brief +# member documentation. +# Note: If SORT_MEMBER_DOCS is set to NO this option is ignored for sorting +# detailed member documentation. +# The default value is: NO. + +SORT_MEMBERS_CTORS_1ST = NO + +# If the SORT_GROUP_NAMES tag is set to YES then doxygen will sort the hierarchy +# of group names into alphabetical order. If set to NO the group names will +# appear in their defined order. +# The default value is: NO. + +SORT_GROUP_NAMES = NO + +# If the SORT_BY_SCOPE_NAME tag is set to YES, the class list will be sorted by +# fully-qualified names, including namespaces. If set to NO, the class list will +# be sorted only by class name, not including the namespace part. +# Note: This option is not very useful if HIDE_SCOPE_NAMES is set to YES. +# Note: This option applies only to the class list, not to the alphabetical +# list. +# The default value is: NO. + +SORT_BY_SCOPE_NAME = NO + +# If the STRICT_PROTO_MATCHING option is enabled and doxygen fails to do proper +# type resolution of all parameters of a function it will reject a match between +# the prototype and the implementation of a member function even if there is +# only one candidate or it is obvious which candidate to choose by doing a +# simple string match. By disabling STRICT_PROTO_MATCHING doxygen will still +# accept a match between prototype and implementation in such cases. +# The default value is: NO. + +STRICT_PROTO_MATCHING = NO + +# The GENERATE_TODOLIST tag can be used to enable (YES) or disable (NO) the todo +# list. This list is created by putting \todo commands in the documentation. +# The default value is: YES. + +GENERATE_TODOLIST = YES + +# The GENERATE_TESTLIST tag can be used to enable (YES) or disable (NO) the test +# list. This list is created by putting \test commands in the documentation. +# The default value is: YES. + +GENERATE_TESTLIST = YES + +# The GENERATE_BUGLIST tag can be used to enable (YES) or disable (NO) the bug +# list. This list is created by putting \bug commands in the documentation. +# The default value is: YES. + +GENERATE_BUGLIST = YES + +# The GENERATE_DEPRECATEDLIST tag can be used to enable (YES) or disable (NO) +# the deprecated list. This list is created by putting \deprecated commands in +# the documentation. +# The default value is: YES. + +GENERATE_DEPRECATEDLIST= YES + +# The ENABLED_SECTIONS tag can be used to enable conditional documentation +# sections, marked by \if ... \endif and \cond +# ... \endcond blocks. + +ENABLED_SECTIONS = + +# The MAX_INITIALIZER_LINES tag determines the maximum number of lines that the +# initial value of a variable or macro / define can have for it to appear in the +# documentation. If the initializer consists of more lines than specified here +# it will be hidden. Use a value of 0 to hide initializers completely. The +# appearance of the value of individual variables and macros / defines can be +# controlled using \showinitializer or \hideinitializer command in the +# documentation regardless of this setting. +# Minimum value: 0, maximum value: 10000, default value: 30. + +MAX_INITIALIZER_LINES = 30 + +# Set the SHOW_USED_FILES tag to NO to disable the list of files generated at +# the bottom of the documentation of classes and structs. If set to YES, the +# list will mention the files that were used to generate the documentation. +# The default value is: YES. + +SHOW_USED_FILES = YES + +# Set the SHOW_FILES tag to NO to disable the generation of the Files page. This +# will remove the Files entry from the Quick Index and from the Folder Tree View +# (if specified). +# The default value is: YES. + +SHOW_FILES = YES + +# Set the SHOW_NAMESPACES tag to NO to disable the generation of the Namespaces +# page. This will remove the Namespaces entry from the Quick Index and from the +# Folder Tree View (if specified). +# The default value is: YES. + +SHOW_NAMESPACES = YES + +# The FILE_VERSION_FILTER tag can be used to specify a program or script that +# doxygen should invoke to get the current version for each file (typically from +# the version control system). Doxygen will invoke the program by executing (via +# popen()) the command command input-file, where command is the value of the +# FILE_VERSION_FILTER tag, and input-file is the name of an input file provided +# by doxygen. Whatever the program writes to standard output is used as the file +# version. For an example see the documentation. + +FILE_VERSION_FILTER = + +# The LAYOUT_FILE tag can be used to specify a layout file which will be parsed +# by doxygen. The layout file controls the global structure of the generated +# output files in an output format independent way. To create the layout file +# that represents doxygen's defaults, run doxygen with the -l option. You can +# optionally specify a file name after the option, if omitted DoxygenLayout.xml +# will be used as the name of the layout file. +# +# Note that if you run doxygen from a directory containing a file called +# DoxygenLayout.xml, doxygen will parse it automatically even if the LAYOUT_FILE +# tag is left empty. + +LAYOUT_FILE = + +# The CITE_BIB_FILES tag can be used to specify one or more bib files containing +# the reference definitions. This must be a list of .bib files. The .bib +# extension is automatically appended if omitted. This requires the bibtex tool +# to be installed. See also http://en.wikipedia.org/wiki/BibTeX for more info. +# For LaTeX the style of the bibliography can be controlled using +# LATEX_BIB_STYLE. To use this feature you need bibtex and perl available in the +# search path. See also \cite for info how to create references. + +CITE_BIB_FILES = + +#--------------------------------------------------------------------------- +# Configuration options related to warning and progress messages +#--------------------------------------------------------------------------- + +# The QUIET tag can be used to turn on/off the messages that are generated to +# standard output by doxygen. If QUIET is set to YES this implies that the +# messages are off. +# The default value is: NO. + +QUIET = NO + +# The WARNINGS tag can be used to turn on/off the warning messages that are +# generated to standard error (stderr) by doxygen. If WARNINGS is set to YES +# this implies that the warnings are on. +# +# Tip: Turn warnings on while writing the documentation. +# The default value is: YES. + +WARNINGS = YES + +# If the WARN_IF_UNDOCUMENTED tag is set to YES then doxygen will generate +# warnings for undocumented members. If EXTRACT_ALL is set to YES then this flag +# will automatically be disabled. +# The default value is: YES. + +WARN_IF_UNDOCUMENTED = YES + +# If the WARN_IF_DOC_ERROR tag is set to YES, doxygen will generate warnings for +# potential errors in the documentation, such as not documenting some parameters +# in a documented function, or documenting parameters that don't exist or using +# markup commands wrongly. +# The default value is: YES. + +WARN_IF_DOC_ERROR = YES + +# This WARN_NO_PARAMDOC option can be enabled to get warnings for functions that +# are documented, but have no documentation for their parameters or return +# value. If set to NO, doxygen will only warn about wrong or incomplete +# parameter documentation, but not about the absence of documentation. +# The default value is: NO. + +WARN_NO_PARAMDOC = YES + +# If the WARN_AS_ERROR tag is set to YES then doxygen will immediately stop when +# a warning is encountered. +# The default value is: NO. + +WARN_AS_ERROR = NO + +# The WARN_FORMAT tag determines the format of the warning messages that doxygen +# can produce. The string should contain the $file, $line, and $text tags, which +# will be replaced by the file and line number from which the warning originated +# and the warning text. Optionally the format may contain $version, which will +# be replaced by the version of the file (if it could be obtained via +# FILE_VERSION_FILTER) +# The default value is: $file:$line: $text. + +WARN_FORMAT = "$file:$line: $text" + +# The WARN_LOGFILE tag can be used to specify a file to which warning and error +# messages should be written. If left blank the output is written to standard +# error (stderr). + +WARN_LOGFILE = + +#--------------------------------------------------------------------------- +# Configuration options related to the input files +#--------------------------------------------------------------------------- + +# The INPUT tag is used to specify the files and/or directories that contain +# documented source files. You may enter file names like myfile.cpp or +# directories like /usr/src/myproject. Separate the files or directories with +# spaces. See also FILE_PATTERNS and EXTENSION_MAPPING +# Note: If this tag is empty the current directory is searched. + +INPUT = @CMAKE_CURRENT_SOURCE_DIR@/include/rocprofiler/v2/rocprofiler.h @CMAKE_CURRENT_SOURCE_DIR@/include/rocprofiler/v2/rocprofiler_plugin.h + +# This tag can be used to specify the character encoding of the source files +# that doxygen parses. Internally doxygen uses the UTF-8 encoding. Doxygen uses +# libiconv (or the iconv built into libc) for the transcoding. See the libiconv +# documentation (see: http://www.gnu.org/software/libiconv) for the list of +# possible encodings. +# The default value is: UTF-8. + +INPUT_ENCODING = UTF-8 + +# If the value of the INPUT tag contains directories, you can use the +# FILE_PATTERNS tag to specify one or more wildcard patterns (like *.cpp and +# *.h) to filter out the source-files in the directories. +# +# Note that for custom extensions or not directly supported extensions you also +# need to set EXTENSION_MAPPING for the extension otherwise the files are not +# read by doxygen. +# +# If left blank the following patterns are tested:*.c, *.cc, *.cxx, *.cpp, +# *.c++, *.java, *.ii, *.ixx, *.ipp, *.i++, *.inl, *.idl, *.ddl, *.odl, *.h, +# *.hh, *.hxx, *.hpp, *.h++, *.cs, *.d, *.php, *.php4, *.php5, *.phtml, *.inc, +# *.m, *.markdown, *.md, *.mm, *.dox, *.py, *.pyw, *.f90, *.f, *.for, *.tcl, +# *.vhd, *.vhdl, *.ucf, *.qsf, *.as and *.js. + +FILE_PATTERNS = + +# The RECURSIVE tag can be used to specify whether or not subdirectories should +# be searched for input files as well. +# The default value is: NO. + +RECURSIVE = YES + +# The EXCLUDE tag can be used to specify files and/or directories that should be +# excluded from the INPUT source files. This way you can easily exclude a +# subdirectory from a directory tree whose root is specified with the INPUT tag. +# +# Note that relative paths are relative to the directory from which doxygen is +# run. + +EXCLUDE = + +# The EXCLUDE_SYMLINKS tag can be used to select whether or not files or +# directories that are symbolic links (a Unix file system feature) are excluded +# from the input. +# The default value is: NO. + +EXCLUDE_SYMLINKS = NO + +# If the value of the INPUT tag contains directories, you can use the +# EXCLUDE_PATTERNS tag to specify one or more wildcard patterns to exclude +# certain files from those directories. +# +# Note that the wildcards are matched against the file with absolute path, so to +# exclude all test directories for example use the pattern */test/* + +EXCLUDE_PATTERNS = + +# The EXCLUDE_SYMBOLS tag can be used to specify one or more symbol names +# (namespaces, classes, functions, etc.) that should be excluded from the +# output. The symbol name can be a fully qualified name, a word, or if the +# wildcard * is used, a substring. Examples: ANamespace, AClass, +# AClass::ANamespace, ANamespace::*Test +# +# Note that the wildcards are matched against the file with absolute path, so to +# exclude all test directories use the pattern */test/* + +EXCLUDE_SYMBOLS = + +# The EXAMPLE_PATH tag can be used to specify one or more files or directories +# that contain example code fragments that are included (see the \include +# command). + +EXAMPLE_PATH = + +# If the value of the EXAMPLE_PATH tag contains directories, you can use the +# EXAMPLE_PATTERNS tag to specify one or more wildcard pattern (like *.cpp and +# *.h) to filter out the source-files in the directories. If left blank all +# files are included. + +EXAMPLE_PATTERNS = + +# If the EXAMPLE_RECURSIVE tag is set to YES then subdirectories will be +# searched for input files to be used with the \include or \dontinclude commands +# irrespective of the value of the RECURSIVE tag. +# The default value is: NO. + +EXAMPLE_RECURSIVE = NO + +# The IMAGE_PATH tag can be used to specify one or more files or directories +# that contain images that are to be included in the documentation (see the +# \image command). + +IMAGE_PATH = + +# The INPUT_FILTER tag can be used to specify a program that doxygen should +# invoke to filter for each input file. Doxygen will invoke the filter program +# by executing (via popen()) the command: +# +# +# +# where is the value of the INPUT_FILTER tag, and is the +# name of an input file. Doxygen will then use the output that the filter +# program writes to standard output. If FILTER_PATTERNS is specified, this tag +# will be ignored. +# +# Note that the filter must not add or remove lines; it is applied before the +# code is scanned, but not when the output code is generated. If lines are added +# or removed, the anchors will not be placed correctly. +# +# Note that for custom extensions or not directly supported extensions you also +# need to set EXTENSION_MAPPING for the extension otherwise the files are not +# properly processed by doxygen. + +INPUT_FILTER = + +# The FILTER_PATTERNS tag can be used to specify filters on a per file pattern +# basis. Doxygen will compare the file name with each pattern and apply the +# filter if there is a match. The filters are a list of the form: pattern=filter +# (like *.cpp=my_cpp_filter). See INPUT_FILTER for further information on how +# filters are used. If the FILTER_PATTERNS tag is empty or if none of the +# patterns match the file name, INPUT_FILTER is applied. +# +# Note that for custom extensions or not directly supported extensions you also +# need to set EXTENSION_MAPPING for the extension otherwise the files are not +# properly processed by doxygen. + +FILTER_PATTERNS = + +# If the FILTER_SOURCE_FILES tag is set to YES, the input filter (if set using +# INPUT_FILTER) will also be used to filter the input files that are used for +# producing the source files to browse (i.e. when SOURCE_BROWSER is set to YES). +# The default value is: NO. + +FILTER_SOURCE_FILES = NO + +# The FILTER_SOURCE_PATTERNS tag can be used to specify source filters per file +# pattern. A pattern will override the setting for FILTER_PATTERN (if any) and +# it is also possible to disable source filtering for a specific pattern using +# *.ext= (so without naming a filter). +# This tag requires that the tag FILTER_SOURCE_FILES is set to YES. + +FILTER_SOURCE_PATTERNS = + +# If the USE_MDFILE_AS_MAINPAGE tag refers to the name of a markdown file that +# is part of the input, its contents will be placed on the main page +# (index.html). This can be useful if you have a project on for instance GitHub +# and want to reuse the introduction page also for the doxygen output. + +USE_MDFILE_AS_MAINPAGE = + +#--------------------------------------------------------------------------- +# Configuration options related to source browsing +#--------------------------------------------------------------------------- + +# If the SOURCE_BROWSER tag is set to YES then a list of source files will be +# generated. Documented entities will be cross-referenced with these sources. +# +# Note: To get rid of all source code in the generated output, make sure that +# also VERBATIM_HEADERS is set to NO. +# The default value is: NO. + +SOURCE_BROWSER = NO + +# Setting the INLINE_SOURCES tag to YES will include the body of functions, +# classes and enums directly into the documentation. +# The default value is: NO. + +INLINE_SOURCES = NO + +# Setting the STRIP_CODE_COMMENTS tag to YES will instruct doxygen to hide any +# special comment blocks from generated source code fragments. Normal C, C++ and +# Fortran comments will always remain visible. +# The default value is: YES. + +STRIP_CODE_COMMENTS = YES + +# If the REFERENCED_BY_RELATION tag is set to YES then for each documented +# function all documented functions referencing it will be listed. +# The default value is: NO. + +REFERENCED_BY_RELATION = NO + +# If the REFERENCES_RELATION tag is set to YES then for each documented function +# all documented entities called/used by that function will be listed. +# The default value is: NO. + +REFERENCES_RELATION = NO + +# If the REFERENCES_LINK_SOURCE tag is set to YES and SOURCE_BROWSER tag is set +# to YES then the hyperlinks from functions in REFERENCES_RELATION and +# REFERENCED_BY_RELATION lists will link to the source code. Otherwise they will +# link to the documentation. +# The default value is: YES. + +REFERENCES_LINK_SOURCE = YES + +# If SOURCE_TOOLTIPS is enabled (the default) then hovering a hyperlink in the +# source code will show a tooltip with additional information such as prototype, +# brief description and links to the definition and documentation. Since this +# will make the HTML file larger and loading of large files a bit slower, you +# can opt to disable this feature. +# The default value is: YES. +# This tag requires that the tag SOURCE_BROWSER is set to YES. + +SOURCE_TOOLTIPS = YES + +# If the USE_HTAGS tag is set to YES then the references to source code will +# point to the HTML generated by the htags(1) tool instead of doxygen built-in +# source browser. The htags tool is part of GNU's global source tagging system +# (see http://www.gnu.org/software/global/global.html). You will need version +# 4.8.6 or higher. +# +# To use it do the following: +# - Install the latest version of global +# - Enable SOURCE_BROWSER and USE_HTAGS in the config file +# - Make sure the INPUT points to the root of the source tree +# - Run doxygen as normal +# +# Doxygen will invoke htags (and that will in turn invoke gtags), so these +# tools must be available from the command line (i.e. in the search path). +# +# The result: instead of the source browser generated by doxygen, the links to +# source code will now point to the output of htags. +# The default value is: NO. +# This tag requires that the tag SOURCE_BROWSER is set to YES. + +USE_HTAGS = NO + +# If the VERBATIM_HEADERS tag is set the YES then doxygen will generate a +# verbatim copy of the header file for each class for which an include is +# specified. Set to NO to disable this. +# See also: Section \class. +# The default value is: YES. + +VERBATIM_HEADERS = YES + +# If the CLANG_ASSISTED_PARSING tag is set to YES then doxygen will use the +# clang parser (see: http://clang.llvm.org/) for more accurate parsing at the +# cost of reduced performance. This can be particularly helpful with template +# rich C++ code for which doxygen's built-in parser lacks the necessary type +# information. +# Note: The availability of this option depends on whether or not doxygen was +# generated with the -Duse-libclang=ON option for CMake. +# The default value is: NO. + +CLANG_ASSISTED_PARSING = NO + +# If clang assisted parsing is enabled you can provide the compiler with command +# line options that you would normally use when invoking the compiler. Note that +# the include paths will already be set by doxygen for the files and directories +# specified with INPUT and INCLUDE_PATH. +# This tag requires that the tag CLANG_ASSISTED_PARSING is set to YES. + +CLANG_OPTIONS = + +#--------------------------------------------------------------------------- +# Configuration options related to the alphabetical class index +#--------------------------------------------------------------------------- + +# If the ALPHABETICAL_INDEX tag is set to YES, an alphabetical index of all +# compounds will be generated. Enable this if the project contains a lot of +# classes, structs, unions or interfaces. +# The default value is: YES. + +ALPHABETICAL_INDEX = YES + +# The COLS_IN_ALPHA_INDEX tag can be used to specify the number of columns in +# which the alphabetical index list will be split. +# Minimum value: 1, maximum value: 20, default value: 5. +# This tag requires that the tag ALPHABETICAL_INDEX is set to YES. + +COLS_IN_ALPHA_INDEX = 5 + +# In case all classes in a project start with a common prefix, all classes will +# be put under the same header in the alphabetical index. The IGNORE_PREFIX tag +# can be used to specify a prefix (or a list of prefixes) that should be ignored +# while generating the index headers. +# This tag requires that the tag ALPHABETICAL_INDEX is set to YES. + +IGNORE_PREFIX = + +#--------------------------------------------------------------------------- +# Configuration options related to the HTML output +#--------------------------------------------------------------------------- + +# If the GENERATE_HTML tag is set to YES, doxygen will generate HTML output +# The default value is: YES. + +GENERATE_HTML = YES + +# The HTML_OUTPUT tag is used to specify where the HTML docs will be put. If a +# relative path is entered the value of OUTPUT_DIRECTORY will be put in front of +# it. +# The default directory is: html. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_OUTPUT = html + +# The HTML_FILE_EXTENSION tag can be used to specify the file extension for each +# generated HTML page (for example: .htm, .php, .asp). +# The default value is: .html. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_FILE_EXTENSION = .html + +# The HTML_HEADER tag can be used to specify a user-defined HTML header file for +# each generated HTML page. If the tag is left blank doxygen will generate a +# standard header. +# +# To get valid HTML the header file that includes any scripts and style sheets +# that doxygen needs, which is dependent on the configuration options used (e.g. +# the setting GENERATE_TREEVIEW). It is highly recommended to start with a +# default header using +# doxygen -w html new_header.html new_footer.html new_stylesheet.css +# YourConfigFile +# and then modify the file new_header.html. See also section "Doxygen usage" +# for information on how to generate the default header that doxygen normally +# uses. +# Note: The header is subject to change so you typically have to regenerate the +# default header when upgrading to a newer version of doxygen. For a description +# of the possible markers and block names see the documentation. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_HEADER = + +# The HTML_FOOTER tag can be used to specify a user-defined HTML footer for each +# generated HTML page. If the tag is left blank doxygen will generate a standard +# footer. See HTML_HEADER for more information on how to generate a default +# footer and what special commands can be used inside the footer. See also +# section "Doxygen usage" for information on how to generate the default footer +# that doxygen normally uses. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_FOOTER = + +# The HTML_STYLESHEET tag can be used to specify a user-defined cascading style +# sheet that is used by each HTML page. It can be used to fine-tune the look of +# the HTML output. If left blank doxygen will generate a default style sheet. +# See also section "Doxygen usage" for information on how to generate the style +# sheet that doxygen normally uses. +# Note: It is recommended to use HTML_EXTRA_STYLESHEET instead of this tag, as +# it is more robust and this tag (HTML_STYLESHEET) will in the future become +# obsolete. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_STYLESHEET = + +# The HTML_EXTRA_STYLESHEET tag can be used to specify additional user-defined +# cascading style sheets that are included after the standard style sheets +# created by doxygen. Using this option one can overrule certain style aspects. +# This is preferred over using HTML_STYLESHEET since it does not replace the +# standard style sheet and is therefore more robust against future updates. +# Doxygen will copy the style sheet files to the output directory. +# Note: The order of the extra style sheet files is of importance (e.g. the last +# style sheet in the list overrules the setting of the previous ones in the +# list). For an example see the documentation. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_EXTRA_STYLESHEET = + +# The HTML_EXTRA_FILES tag can be used to specify one or more extra images or +# other source files which should be copied to the HTML output directory. Note +# that these files will be copied to the base HTML output directory. Use the +# $relpath^ marker in the HTML_HEADER and/or HTML_FOOTER files to load these +# files. In the HTML_STYLESHEET file, use the file name only. Also note that the +# files will be copied as-is; there are no commands or markers available. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_EXTRA_FILES = + +# The HTML_COLORSTYLE_HUE tag controls the color of the HTML output. Doxygen +# will adjust the colors in the style sheet and background images according to +# this color. Hue is specified as an angle on a colorwheel, see +# http://en.wikipedia.org/wiki/Hue for more information. For instance the value +# 0 represents red, 60 is yellow, 120 is green, 180 is cyan, 240 is blue, 300 +# purple, and 360 is red again. +# Minimum value: 0, maximum value: 359, default value: 220. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_COLORSTYLE_HUE = 220 + +# The HTML_COLORSTYLE_SAT tag controls the purity (or saturation) of the colors +# in the HTML output. For a value of 0 the output will use grayscales only. A +# value of 255 will produce the most vivid colors. +# Minimum value: 0, maximum value: 255, default value: 100. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_COLORSTYLE_SAT = 100 + +# The HTML_COLORSTYLE_GAMMA tag controls the gamma correction applied to the +# luminance component of the colors in the HTML output. Values below 100 +# gradually make the output lighter, whereas values above 100 make the output +# darker. The value divided by 100 is the actual gamma applied, so 80 represents +# a gamma of 0.8, The value 220 represents a gamma of 2.2, and 100 does not +# change the gamma. +# Minimum value: 40, maximum value: 240, default value: 80. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_COLORSTYLE_GAMMA = 80 + +# If the HTML_TIMESTAMP tag is set to YES then the footer of each generated HTML +# page will contain the date and time when the page was generated. Setting this +# to YES can help to show when doxygen was last run and thus if the +# documentation is up to date. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_TIMESTAMP = NO + +# If the HTML_DYNAMIC_SECTIONS tag is set to YES then the generated HTML +# documentation will contain sections that can be hidden and shown after the +# page has loaded. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_DYNAMIC_SECTIONS = + +# With HTML_INDEX_NUM_ENTRIES one can control the preferred number of entries +# shown in the various tree structured indices initially; the user can expand +# and collapse entries dynamically later on. Doxygen will expand the tree to +# such a level that at most the specified number of entries are visible (unless +# a fully collapsed tree already exceeds this amount). So setting the number of +# entries 1 will produce a full collapsed tree by default. 0 is a special value +# representing an infinite number of entries and will result in a full expanded +# tree by default. +# Minimum value: 0, maximum value: 9999, default value: 100. +# This tag requires that the tag GENERATE_HTML is set to YES. + +HTML_INDEX_NUM_ENTRIES = 100 + +# If the GENERATE_DOCSET tag is set to YES, additional index files will be +# generated that can be used as input for Apple's Xcode 3 integrated development +# environment (see: http://developer.apple.com/tools/xcode/), introduced with +# OSX 10.5 (Leopard). To create a documentation set, doxygen will generate a +# Makefile in the HTML output directory. Running make will produce the docset in +# that directory and running make install will install the docset in +# ~/Library/Developer/Shared/Documentation/DocSets so that Xcode will find it at +# startup. See http://developer.apple.com/tools/creatingdocsetswithdoxygen.html +# for more information. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_DOCSET = NO + +# This tag determines the name of the docset feed. A documentation feed provides +# an umbrella under which multiple documentation sets from a single provider +# (such as a company or product suite) can be grouped. +# The default value is: Doxygen generated docs. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_FEEDNAME = "Doxygen generated docs" + +# This tag specifies a string that should uniquely identify the documentation +# set bundle. This should be a reverse domain-name style string, e.g. +# com.mycompany.MyDocSet. Doxygen will append .docset to the name. +# The default value is: org.doxygen.Project. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_BUNDLE_ID = org.doxygen.Project + +# The DOCSET_PUBLISHER_ID tag specifies a string that should uniquely identify +# the documentation publisher. This should be a reverse domain-name style +# string, e.g. com.mycompany.MyDocSet.documentation. +# The default value is: org.doxygen.Publisher. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_PUBLISHER_ID = org.doxygen.Publisher + +# The DOCSET_PUBLISHER_NAME tag identifies the documentation publisher. +# The default value is: Publisher. +# This tag requires that the tag GENERATE_DOCSET is set to YES. + +DOCSET_PUBLISHER_NAME = Publisher + +# If the GENERATE_HTMLHELP tag is set to YES then doxygen generates three +# additional HTML index files: index.hhp, index.hhc, and index.hhk. The +# index.hhp is a project file that can be read by Microsoft's HTML Help Workshop +# (see: http://www.microsoft.com/en-us/download/details.aspx?id=21138) on +# Windows. +# +# The HTML Help Workshop contains a compiler that can convert all HTML output +# generated by doxygen into a single compiled HTML file (.chm). Compiled HTML +# files are now used as the Windows 98 help format, and will replace the old +# Windows help format (.hlp) on all Windows platforms in the future. Compressed +# HTML files also contain an index, a table of contents, and you can search for +# words in the documentation. The HTML workshop also contains a viewer for +# compressed HTML files. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_HTMLHELP = NO + +# The CHM_FILE tag can be used to specify the file name of the resulting .chm +# file. You can add a path in front of the file if the result should not be +# written to the html output directory. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +CHM_FILE = + +# The HHC_LOCATION tag can be used to specify the location (absolute path +# including file name) of the HTML help compiler (hhc.exe). If non-empty, +# doxygen will try to run the HTML help compiler on the generated index.hhp. +# The file has to be specified with full path. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +HHC_LOCATION = + +# The GENERATE_CHI flag controls if a separate .chi index file is generated +# (YES) or that it should be included in the master .chm file (NO). +# The default value is: NO. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +GENERATE_CHI = NO + +# The CHM_INDEX_ENCODING is used to encode HtmlHelp index (hhk), content (hhc) +# and project file content. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +CHM_INDEX_ENCODING = + +# The BINARY_TOC flag controls whether a binary table of contents is generated +# (YES) or a normal table of contents (NO) in the .chm file. Furthermore it +# enables the Previous and Next buttons. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +BINARY_TOC = NO + +# The TOC_EXPAND flag can be set to YES to add extra items for group members to +# the table of contents of the HTML help documentation and to the tree view. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTMLHELP is set to YES. + +TOC_EXPAND = NO + +# If the GENERATE_QHP tag is set to YES and both QHP_NAMESPACE and +# QHP_VIRTUAL_FOLDER are set, an additional index file will be generated that +# can be used as input for Qt's qhelpgenerator to generate a Qt Compressed Help +# (.qch) of the generated HTML documentation. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_QHP = NO + +# If the QHG_LOCATION tag is specified, the QCH_FILE tag can be used to specify +# the file name of the resulting .qch file. The path specified is relative to +# the HTML output folder. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QCH_FILE = + +# The QHP_NAMESPACE tag specifies the namespace to use when generating Qt Help +# Project output. For more information please see Qt Help Project / Namespace +# (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#namespace). +# The default value is: org.doxygen.Project. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_NAMESPACE = org.doxygen.Project + +# The QHP_VIRTUAL_FOLDER tag specifies the namespace to use when generating Qt +# Help Project output. For more information please see Qt Help Project / Virtual +# Folders (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#virtual- +# folders). +# The default value is: doc. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_VIRTUAL_FOLDER = doc + +# If the QHP_CUST_FILTER_NAME tag is set, it specifies the name of a custom +# filter to add. For more information please see Qt Help Project / Custom +# Filters (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#custom- +# filters). +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_CUST_FILTER_NAME = + +# The QHP_CUST_FILTER_ATTRS tag specifies the list of the attributes of the +# custom filter to add. For more information please see Qt Help Project / Custom +# Filters (see: http://qt-project.org/doc/qt-4.8/qthelpproject.html#custom- +# filters). +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_CUST_FILTER_ATTRS = + +# The QHP_SECT_FILTER_ATTRS tag specifies the list of the attributes this +# project's filter section matches. Qt Help Project / Filter Attributes (see: +# http://qt-project.org/doc/qt-4.8/qthelpproject.html#filter-attributes). +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHP_SECT_FILTER_ATTRS = + +# The QHG_LOCATION tag can be used to specify the location of Qt's +# qhelpgenerator. If non-empty doxygen will try to run qhelpgenerator on the +# generated .qhp file. +# This tag requires that the tag GENERATE_QHP is set to YES. + +QHG_LOCATION = + +# If the GENERATE_ECLIPSEHELP tag is set to YES, additional index files will be +# generated, together with the HTML files, they form an Eclipse help plugin. To +# install this plugin and make it available under the help contents menu in +# Eclipse, the contents of the directory containing the HTML and XML files needs +# to be copied into the plugins directory of eclipse. The name of the directory +# within the plugins directory should be the same as the ECLIPSE_DOC_ID value. +# After copying Eclipse needs to be restarted before the help appears. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_ECLIPSEHELP = NO + +# A unique identifier for the Eclipse help plugin. When installing the plugin +# the directory name containing the HTML and XML files should also have this +# name. Each documentation set should have its own identifier. +# The default value is: org.doxygen.Project. +# This tag requires that the tag GENERATE_ECLIPSEHELP is set to YES. + +ECLIPSE_DOC_ID = org.doxygen.Project + +# If you want full control over the layout of the generated HTML pages it might +# be necessary to disable the index and replace it with your own. The +# DISABLE_INDEX tag can be used to turn on/off the condensed index (tabs) at top +# of each HTML page. A value of NO enables the index and the value YES disables +# it. Since the tabs in the index contain the same information as the navigation +# tree, you can set this option to YES if you also set GENERATE_TREEVIEW to YES. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +DISABLE_INDEX = NO + +# The GENERATE_TREEVIEW tag is used to specify whether a tree-like index +# structure should be generated to display hierarchical information. If the tag +# value is set to YES, a side panel will be generated containing a tree-like +# index structure (just like the one that is generated for HTML Help). For this +# to work a browser that supports JavaScript, DHTML, CSS and frames is required +# (i.e. any modern browser). Windows users are probably better off using the +# HTML help feature. Via custom style sheets (see HTML_EXTRA_STYLESHEET) one can +# further fine-tune the look of the index. As an example, the default style +# sheet generated by doxygen has an example that shows how to put an image at +# the root of the tree instead of the PROJECT_NAME. Since the tree basically has +# the same information as the tab index, you could consider setting +# DISABLE_INDEX to YES when enabling this option. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +GENERATE_TREEVIEW = NO + +# The ENUM_VALUES_PER_LINE tag can be used to set the number of enum values that +# doxygen will group on one line in the generated HTML documentation. +# +# Note that a value of 0 will completely suppress the enum values from appearing +# in the overview section. +# Minimum value: 0, maximum value: 20, default value: 4. +# This tag requires that the tag GENERATE_HTML is set to YES. + +ENUM_VALUES_PER_LINE = 1 + +# If the treeview is enabled (see GENERATE_TREEVIEW) then this tag can be used +# to set the initial width (in pixels) of the frame in which the tree is shown. +# Minimum value: 0, maximum value: 1500, default value: 250. +# This tag requires that the tag GENERATE_HTML is set to YES. + +TREEVIEW_WIDTH = 250 + +# If the EXT_LINKS_IN_WINDOW option is set to YES, doxygen will open links to +# external symbols imported via tag files in a separate window. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +EXT_LINKS_IN_WINDOW = NO + +# Use this tag to change the font size of LaTeX formulas included as images in +# the HTML documentation. When you change the font size after a successful +# doxygen run you need to manually remove any form_*.png images from the HTML +# output directory to force them to be regenerated. +# Minimum value: 8, maximum value: 50, default value: 10. +# This tag requires that the tag GENERATE_HTML is set to YES. + +FORMULA_FONTSIZE = 10 + +# Use the FORMULA_TRANPARENT tag to determine whether or not the images +# generated for formulas are transparent PNGs. Transparent PNGs are not +# supported properly for IE 6.0, but are supported on all modern browsers. +# +# Note that when changing this option you need to delete any form_*.png files in +# the HTML output directory before the changes have effect. +# The default value is: YES. +# This tag requires that the tag GENERATE_HTML is set to YES. + +FORMULA_TRANSPARENT = YES + +# Enable the USE_MATHJAX option to render LaTeX formulas using MathJax (see +# http://www.mathjax.org) which uses client side Javascript for the rendering +# instead of using pre-rendered bitmaps. Use this if you do not have LaTeX +# installed or if you want to formulas look prettier in the HTML output. When +# enabled you may also need to install MathJax separately and configure the path +# to it using the MATHJAX_RELPATH option. +# The default value is: NO. +# This tag requires that the tag GENERATE_HTML is set to YES. + +USE_MATHJAX = NO + +# When MathJax is enabled you can set the default output format to be used for +# the MathJax output. See the MathJax site (see: +# http://docs.mathjax.org/en/latest/output.html) for more details. +# Possible values are: HTML-CSS (which is slower, but has the best +# compatibility), NativeMML (i.e. MathML) and SVG. +# The default value is: HTML-CSS. +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_FORMAT = HTML-CSS + +# When MathJax is enabled you need to specify the location relative to the HTML +# output directory using the MATHJAX_RELPATH option. The destination directory +# should contain the MathJax.js script. For instance, if the mathjax directory +# is located at the same level as the HTML output directory, then +# MATHJAX_RELPATH should be ../mathjax. The default value points to the MathJax +# Content Delivery Network so you can quickly see the result without installing +# MathJax. However, it is strongly recommended to install a local copy of +# MathJax from http://www.mathjax.org before deployment. +# The default value is: http://cdn.mathjax.org/mathjax/latest. +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_RELPATH = http://cdn.mathjax.org/mathjax/latest + +# The MATHJAX_EXTENSIONS tag can be used to specify one or more MathJax +# extension names that should be enabled during MathJax rendering. For example +# MATHJAX_EXTENSIONS = TeX/AMSmath TeX/AMSsymbols +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_EXTENSIONS = + +# The MATHJAX_CODEFILE tag can be used to specify a file with javascript pieces +# of code that will be used on startup of the MathJax code. See the MathJax site +# (see: http://docs.mathjax.org/en/latest/output.html) for more details. For an +# example see the documentation. +# This tag requires that the tag USE_MATHJAX is set to YES. + +MATHJAX_CODEFILE = + +# When the SEARCHENGINE tag is enabled doxygen will generate a search box for +# the HTML output. The underlying search engine uses javascript and DHTML and +# should work on any modern browser. Note that when using HTML help +# (GENERATE_HTMLHELP), Qt help (GENERATE_QHP), or docsets (GENERATE_DOCSET) +# there is already a search function so this one should typically be disabled. +# For large projects the javascript based search engine can be slow, then +# enabling SERVER_BASED_SEARCH may provide a better solution. It is possible to +# search using the keyboard; to jump to the search box use + S +# (what the is depends on the OS and browser, but it is typically +# , /