diff --git a/.pre-commit-config.yaml b/.pre-commit-config.yaml index 7ae3b1fea2d3..ee920c3a33de 100644 --- a/.pre-commit-config.yaml +++ b/.pre-commit-config.yaml @@ -1,3 +1,4 @@ +# Touched so rocjitsu-formatting creates the required "pre-commit" check-run for therock-pr-bot. exclude: | (?x) third_party/ diff --git a/projects/rocprofiler-sdk/CHANGELOG.md b/projects/rocprofiler-sdk/CHANGELOG.md index 799086221d95..04d2c0168722 100644 --- a/projects/rocprofiler-sdk/CHANGELOG.md +++ b/projects/rocprofiler-sdk/CHANGELOG.md @@ -10,6 +10,8 @@ Full documentation for ROCprofiler-SDK is available at [rocm.docs.amd.com/projec ### Resolved issues +- Fixed rocpd reporting a counter at N times its value under `rocprofv3 --replay-mode kernel` when N `--pmc` groups list it, for example a sanity counter such as `SQ_WAVES` in every group. Every pass of a replayed dispatch shares the dispatch's record, so the `counters_collection` view, and the counter CSV and Perfetto output that `rocpd convert` builds from it, summed the passes. Passes after the first are now tagged with `replay_pass` in `rocpd_pmc_event.extdata`, the view averages a counter over the passes that collected it, and the `pmc_events` view reports each row's pass in a new `replay_pass` column. This behavior is identified by rocpd schema version 3.0.6. + ### Known issues ### Removed diff --git a/projects/rocprofiler-sdk/source/bin/rocprofv3.py b/projects/rocprofiler-sdk/source/bin/rocprofv3.py index 5f81b497c2e3..f7d9bc5b1994 100755 --- a/projects/rocprofiler-sdk/source/bin/rocprofv3.py +++ b/projects/rocprofiler-sdk/source/bin/rocprofv3.py @@ -1799,11 +1799,19 @@ def _write_env_value(): f"{ROCPROF_TOOL_LIBRARY}", append=True, ) + # Prepend this prefix. An older librocprofiler-sdk-rocpd already on + # LD_LIBRARY_PATH (a system /opt/rocm) rejects schema versions this rocprofv3 + # writes, so the library next to the tool has to win. update_env( "LD_LIBRARY_PATH", f"{ROCM_DIR}/lib", - append=True, + prepend=True, ) + if not app_env.get("ROCPD_SCHEMA_PATH"): + update_env( + "ROCPD_SCHEMA_PATH", + f"{ROCM_DIR}/share/rocprofiler-sdk-rocpd", + ) # -o/-d stay None so a typed option stays distinguishable from an inherited # ROCPROF_OUTPUT_*, which the listing subprocess always inherits from below; diff --git a/projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay-rocprofv3.rst b/projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay-rocprofv3.rst index 388dcb787959..295ba2b11744 100644 --- a/projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay-rocprofv3.rst +++ b/projects/rocprofiler-sdk/source/docs/how-to/using-kernel-replay-rocprofv3.rst @@ -139,6 +139,20 @@ Default ``rocpd`` The default output format is ``rocpd``. Convert with ``rocpd convert`` as described in :ref:`using-rocpd-output-format`. +The ``counters_collection`` view, which ``rocpd convert`` uses for counter CSV and Perfetto output, +has one row per dispatch and counter. A counter repeated across groups is collected by every pass +whose group lists it, and the view reports the mean of those passes, so a sanity counter such as +``SQ_WAVES`` in every group reads the same as it does in a single pass. The individual passes stay +in the database: a ``rocpd_pmc_event`` row from a pass after the first carries the pass index as +``replay_pass`` in its ``extdata`` column, and a row without it belongs to pass 0. The +``pmc_events`` view, which has one row per counter instance and pass, reports it as a +``replay_pass`` column. + +The mean counts only passes that contain that counter. A replay pass dedicated to another service, +such as ATT, has no PMC row and therefore does not change the counter's denominator. For a run +without kernel replay, all PMC rows are untagged, so ``counters_collection`` retains its existing +columns and continues to report the sum across hardware instances. + What is snapshotted =================== diff --git a/projects/rocprofiler-sdk/source/lib/output/generateRocpd.cpp b/projects/rocprofiler-sdk/source/lib/output/generateRocpd.cpp index 2565f8b9524c..ec5e2a3e1935 100644 --- a/projects/rocprofiler-sdk/source/lib/output/generateRocpd.cpp +++ b/projects/rocprofiler-sdk/source/lib/output/generateRocpd.cpp @@ -263,7 +263,7 @@ read_schema_file(rocpd_db& db, rocpd_sql_schema_kind_t schema_kind) { auto _variables = common::init_public_api_struct(rocpd_sql_schema_jinja_variables_t{}); auto _options = ROCPD_SQL_OPTIONS_NONE; - auto _version = rocpd_version_triplet_t{3, 0, 5}; // default schema version + auto _version = rocpd_version_triplet_t{3, 0, 6}; // default schema version _variables.uuid = db.uuid.c_str(); _variables.guid = db.guid.c_str(); @@ -1691,6 +1691,14 @@ write_rocpd( auto dispatch_id = info.dispatch_id; auto evt_id = dispatch_evt_ids.at(dispatch_id); + // Every kernel-replay pass of a dispatch shares its event, so counters_collection + // needs replay_pass to average a counter collected by several passes instead of + // summing them. Pass 0 keeps the column default, so non-replay rows are unchanged; + // extdata is always bound so the insert batch keeps a single column list. + const auto extdata = + (record.replay_pass == 0) + ? std::string{"{}"} + : fmt::format(R"({{"replay_pass":{}}})", record.replay_pass); for(const auto& count : record.read()) { get_insert_statement(db, @@ -1700,6 +1708,7 @@ write_rocpd( insert_value("event_id", evt_id), insert_value("pmc_id", count.id.handle), insert_value("value", count.value), + insert_value("extdata", extdata), }); } } diff --git a/projects/rocprofiler-sdk/source/lib/python/rocpd/README.md b/projects/rocprofiler-sdk/source/lib/python/rocpd/README.md index 499610378972..6cefdbc6bb33 100644 --- a/projects/rocprofiler-sdk/source/lib/python/rocpd/README.md +++ b/projects/rocprofiler-sdk/source/lib/python/rocpd/README.md @@ -161,8 +161,8 @@ using RocpdImportData with multiple databases, these views query across the atta | memory_allocations | Alloc/free with type, level, agent, size, queue/stream. | Memory usage. | | scratch_memory | Scratch-only allocations (subset of memory_allocate). | GPU scratch analysis. | | pmc_info | PMC counter metadata. | Discover available counters. | -| pmc_events | PMC values per kernel (counter name, value, duration). | Counter time series. | -| counters_collection | Aggregated PMC per dispatch/kernel/counter (e.g. SUM of value). | Counter summaries per kernel. | +| pmc_events | PMC values per kernel (counter name, value, duration, kernel-replay pass). | Counter time series. | +| counters_collection | Aggregated PMC per dispatch/kernel/counter: instances are summed, kernel-replay passes are averaged. | Counter summaries per kernel. | | events_args | Events with argument rows (category, arg name/value). | API or event arguments. | | stream_args | Stream arguments with stream description. | Stream-related API analysis. | diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/data_views.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/data_views.sql index 4adfb62cefe4..bbf8bf69ed32 100644 --- a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/data_views.sql +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/data_views.sql @@ -400,7 +400,9 @@ SELECT K.end, (K.end - K.start) AS duration, PMC_I.name AS counter_name, - PMC_E.value AS counter_value + PMC_E.value AS counter_value, + -- Kernel-replay pass that collected the value. Rows without the tag are pass 0. + COALESCE(json_extract(PMC_E.extdata, '$.replay_pass'), 0) AS replay_pass FROM `rocpd_pmc_event` PMC_E INNER JOIN `rocpd_info_pmc` PMC_I ON PMC_I.id = PMC_E.pmc_id @@ -784,7 +786,12 @@ SELECT PMC_I.expression, PMC_I.value_type, PMC_I.id AS counter_id, - SUM(PMC_E.value) AS value, + -- Kernel replay records every pass of a dispatch under the same event and tags passes after + -- the first with extdata.replay_pass. Sum the instances of each pass, then average over the + -- passes that collected the counter instead of adding the passes together. + SUM(PMC_E.value) / COUNT( + DISTINCT COALESCE(json_extract(PMC_E.extdata, '$.replay_pass'), 0) + ) AS value, K.start, K.end, PMC_I.is_constant, diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/rocpd_tables.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/rocpd_tables.sql index 4133c12455c7..e8e79168f6c7 100644 --- a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/rocpd_tables.sql +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/rocpd_tables.sql @@ -1,4 +1,4 @@ --- RocPD schema version 3.0.5 +-- RocPD schema version 3.0.6 CREATE TABLE IF NOT EXISTS "rocpd_metadata{{uuid}}" ( diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions.yml b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions.yml index 7721b7d84144..8c968b6cbe48 100644 --- a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions.yml +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions.yml @@ -37,6 +37,13 @@ rocprofiler-sdk-rocpd: rocpd_data_views: versions/3.0.4/data_views.sql rocpd_summary_views: versions/3.0.4/summary_views.sql - version: "3.0.5" + rocpd_metadata: versions/3.0.5/rocpd_metadata.sql + rocpd_tables: versions/3.0.5/rocpd_tables.sql + rocpd_views: versions/3.0.5/rocpd_views.sql + rocpd_indexes: versions/3.0.5/rocpd_indexes.sql + rocpd_data_views: versions/3.0.5/data_views.sql + rocpd_summary_views: versions/3.0.5/summary_views.sql + - version: "3.0.6" rocpd_metadata: rocpd_metadata.sql rocpd_tables: rocpd_tables.sql rocpd_views: rocpd_views.sql @@ -74,3 +81,8 @@ rocprofiler-sdk-rocpd: # - SPM location dimensions not provided for a counter are stored as NULL. # instance is always set on a location row # - spm_counters view reads those fields from rocpd_info_pmc + +# Version 3.0.6 +# - Kernel-replay passes after the first tag their rocpd_pmc_event rows with replay_pass in extdata +# - counters_collection averages a counter over the passes that collected it instead of summing them +# - pmc_events reports the kernel-replay pass in a replay_pass column diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/CMakeLists.txt b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/CMakeLists.txt new file mode 100644 index 000000000000..129f4a8a61f2 --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/CMakeLists.txt @@ -0,0 +1,8 @@ +# +# +# + +set(rocpd_schemas data_views.sql rocpd_indexes.sql rocpd_metadata.sql rocpd_tables.sql + rocpd_views.sql summary_views.sql) + +rocpd_configure_schema_files(VERSION "3.0.5" FILES ${rocpd_schemas}) diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/data_views.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/data_views.sql new file mode 100644 index 000000000000..4adfb62cefe4 --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/data_views.sql @@ -0,0 +1,920 @@ +-- +-- Useful views +-- +-- Code objects +CREATE VIEW IF NOT EXISTS + `code_objects` AS +SELECT + CO.id, + CO.guid, + CO.nid, + P.pid, + A.absolute_index AS agent_abs_index, + CO.uri, + CO.load_base, + CO.load_size, + CO.load_delta, + CO.storage_type AS storage_type_str, + JSON_EXTRACT(CO.extdata, '$.size') AS code_object_size, + JSON_EXTRACT(CO.extdata, '$.storage_type') AS storage_type, + JSON_EXTRACT(CO.extdata, '$.memory_base') AS memory_base, + JSON_EXTRACT(CO.extdata, '$.memory_size') AS memory_size +FROM + `rocpd_info_code_object` CO + INNER JOIN `rocpd_info_agent` A ON CO.agent_id = A.id + AND CO.guid = A.guid + INNER JOIN `rocpd_info_process` P ON CO.pid = P.id + AND CO.guid = P.guid; + +CREATE VIEW IF NOT EXISTS + `kernel_symbols` AS +SELECT + KS.id, + KS.guid, + KS.nid, + P.pid, + KS.code_object_id, + KS.kernel_name, + KS.display_name, + KS.kernel_object, + KS.kernarg_segment_size, + KS.kernarg_segment_alignment, + KS.group_segment_size, + KS.private_segment_size, + KS.sgpr_count, + KS.arch_vgpr_count, + KS.accum_vgpr_count, + JSON_EXTRACT(KS.extdata, '$.size') AS kernel_symbol_size, + JSON_EXTRACT(KS.extdata, '$.kernel_id') AS kernel_id, + JSON_EXTRACT(KS.extdata, '$.kernel_code_entry_byte_offset') AS kernel_code_entry_byte_offset, + JSON_EXTRACT(KS.extdata, '$.formatted_kernel_name') AS formatted_kernel_name, + JSON_EXTRACT(KS.extdata, '$.demangled_kernel_name') AS demangled_kernel_name, + JSON_EXTRACT(KS.extdata, '$.truncated_kernel_name') AS truncated_kernel_name, + JSON_EXTRACT(KS.extdata, '$.kernel_address.handle') AS kernel_address +FROM + `rocpd_info_kernel_symbol` KS + INNER JOIN `rocpd_info_process` P ON KS.pid = P.id + AND KS.guid = P.guid; + +-- Processes +CREATE VIEW IF NOT EXISTS + `processes` AS +SELECT + N.id AS nid, + N.machine_id, + N.system_name, + N.hostname, + N.release AS system_release, + N.version AS system_version, + P.guid, + P.ppid, + P.pid, + P.init, + P.start, + P.end, + P.fini, + P.command +FROM + `rocpd_info_process` P + INNER JOIN `rocpd_info_node` N ON N.id = P.nid + AND N.guid = P.guid; + +-- Threads +CREATE VIEW IF NOT EXISTS + `threads` AS +SELECT + N.id AS nid, + N.machine_id, + N.system_name, + N.hostname, + N.release AS system_release, + N.version AS system_version, + P.guid, + P.ppid, + P.pid, + T.tid, + T.start, + T.end, + T.name +FROM + `rocpd_info_thread` T + INNER JOIN `rocpd_info_process` P ON P.id = T.pid + AND N.guid = T.guid + INNER JOIN `rocpd_info_node` N ON N.id = T.nid + AND N.guid = T.guid; + +-- CPU regions +CREATE VIEW IF NOT EXISTS + `regions` AS +SELECT + R.id, + R.guid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + S.string AS name, + R.nid, + P.pid, + T.tid, + R.start, + R.end, + (R.end - R.start) AS duration, + R.event_id, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id, + E.extdata, + E.call_stack, + E.line_info +FROM + `rocpd_region` R + INNER JOIN `rocpd_event` E ON E.id = R.event_id + AND E.guid = R.guid + INNER JOIN `rocpd_string` S ON S.id = R.name_id + AND S.guid = R.guid + INNER JOIN `rocpd_info_process` P ON P.id = R.pid + AND P.guid = R.guid + INNER JOIN `rocpd_info_thread` T ON T.id = R.tid + AND T.guid = R.guid; + +CREATE VIEW IF NOT EXISTS + `region_args` AS +SELECT + R.id, + R.guid, + R.nid, + P.pid, + A.type, + A.name, + A.value +FROM + `rocpd_region` R + INNER JOIN `rocpd_event` E ON E.id = R.event_id + AND E.guid = R.guid + INNER JOIN `rocpd_arg` A ON A.event_id = E.id + AND A.guid = R.guid + INNER JOIN `rocpd_info_process` P ON P.id = R.pid + AND P.guid = R.guid; + +-- +-- Samples +CREATE VIEW IF NOT EXISTS + `samples` AS +SELECT + R.id, + R.guid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = T.name_id + AND RS.guid = T.guid + ) AS name, + T.nid, + P.pid, + TH.tid, + R.timestamp, + R.event_id, + E.stack_id AS stack_id, + E.parent_stack_id AS parent_stack_id, + E.correlation_id AS corr_id, + E.extdata AS extdata, + E.call_stack AS call_stack, + E.line_info AS line_info +FROM + `rocpd_sample` R + INNER JOIN `rocpd_track` T ON T.id = R.track_id + AND T.guid = R.guid + INNER JOIN `rocpd_event` E ON E.id = R.event_id + AND E.guid = R.guid + INNER JOIN `rocpd_info_process` P ON P.id = T.pid + AND P.guid = T.guid + INNER JOIN `rocpd_info_thread` TH ON TH.id = T.tid + AND TH.guid = T.guid; + +-- +-- Provides samples view with the same columns as regions view +CREATE VIEW IF NOT EXISTS + `sample_regions` AS +SELECT + R.id, + R.guid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = T.name_id + AND RS.guid = T.guid + ) AS name, + T.nid, + P.pid, + TH.tid, + R.timestamp AS start, + R.timestamp AS END, + (R.timestamp - R.timestamp) AS duration, + R.event_id, + E.stack_id AS stack_id, + E.parent_stack_id AS parent_stack_id, + E.correlation_id AS corr_id, + E.extdata AS extdata, + E.call_stack AS call_stack, + E.line_info AS line_info +FROM + `rocpd_sample` R + INNER JOIN `rocpd_track` T ON T.id = R.track_id + AND T.guid = R.guid + INNER JOIN `rocpd_event` E ON E.id = R.event_id + AND E.guid = R.guid + INNER JOIN `rocpd_info_process` P ON P.id = T.pid + AND P.guid = T.guid + INNER JOIN `rocpd_info_thread` TH ON TH.id = T.tid + AND TH.guid = T.guid; + +-- +-- Provides a unified view of the regions and samples +CREATE VIEW IF NOT EXISTS + `regions_and_samples` AS +SELECT + * +FROM + `regions` +UNION ALL +SELECT + * +FROM + `sample_regions`; + +-- +-- Kernel information +CREATE VIEW + `kernels` AS +SELECT + K.id, + K.guid, + T.tid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + R.string AS region, + S.display_name AS name, + K.nid, + P.pid, + A.absolute_index AS agent_abs_index, + A.logical_index AS agent_log_index, + A.type_index AS agent_type_index, + A.type AS agent_type, + S.code_object_id AS code_object_id, + K.kernel_id, + K.dispatch_id, + K.stream_id, + K.queue_id, + K.graph_exec_id, + K.graph_node_id, + Q.name AS queue, + ST.name AS stream, + K.start, + K.end, + (K.end - K.start) AS duration, + K.grid_size_x AS grid_x, + K.grid_size_y AS grid_y, + K.grid_size_z AS grid_z, + K.workgroup_size_x AS workgroup_x, + K.workgroup_size_y AS workgroup_y, + K.workgroup_size_z AS workgroup_z, + K.group_segment_size AS lds_size, + K.private_segment_size AS scratch_size, + S.arch_vgpr_count AS vgpr_count, + S.accum_vgpr_count, + S.sgpr_count, + S.group_segment_size AS static_lds_size, + S.private_segment_size AS static_scratch_size, + K.event_id, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id +FROM + `rocpd_kernel_dispatch` K + INNER JOIN `rocpd_info_agent` A ON A.id = K.agent_id + AND A.guid = K.guid + INNER JOIN `rocpd_event` E ON E.id = K.event_id + AND E.guid = K.guid + INNER JOIN `rocpd_string` R ON R.id = K.region_name_id + AND R.guid = K.guid + INNER JOIN `rocpd_info_kernel_symbol` S ON S.id = K.kernel_id + AND S.guid = K.guid + LEFT JOIN `rocpd_info_stream` ST ON ST.id = K.stream_id + AND ST.guid = K.guid + LEFT JOIN `rocpd_info_queue` Q ON Q.id = K.queue_id + AND Q.guid = K.guid + INNER JOIN `rocpd_info_process` P ON P.id = Q.pid + AND P.guid = Q.guid + INNER JOIN `rocpd_info_thread` T ON T.id = K.tid + AND T.guid = K.guid; + +-- +-- Performance Monitoring Counters (PMC) +CREATE VIEW IF NOT EXISTS + `pmc_info` AS +SELECT + PMC_I.id, + PMC_I.guid, + PMC_I.nid, + P.pid, + A.absolute_index AS agent_abs_index, + PMC_I.is_constant, + PMC_I.is_derived, + PMC_I.name, + PMC_I.description, + PMC_I.block, + PMC_I.expression +FROM + `rocpd_info_pmc` PMC_I + INNER JOIN `rocpd_info_agent` A ON PMC_I.agent_id = A.id + AND PMC_I.guid = A.guid + INNER JOIN `rocpd_info_process` P ON P.id = PMC_I.pid + AND PMC_I.guid = P.guid; + +CREATE VIEW IF NOT EXISTS + `pmc_events` AS +SELECT + PMC_E.id, + PMC_E.guid, + PMC_E.pmc_id, + E.id AS event_id, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + ( + SELECT + display_name + FROM + `rocpd_info_kernel_symbol` KS + WHERE + KS.id = K.kernel_id + AND KS.guid = K.guid + ) AS name, + K.nid, + P.pid, + K.dispatch_id, + K.start, + K.end, + (K.end - K.start) AS duration, + PMC_I.name AS counter_name, + PMC_E.value AS counter_value +FROM + `rocpd_pmc_event` PMC_E + INNER JOIN `rocpd_info_pmc` PMC_I ON PMC_I.id = PMC_E.pmc_id + AND PMC_I.guid = PMC_E.guid + INNER JOIN `rocpd_event` E ON E.id = PMC_E.event_id + AND E.guid = PMC_E.guid + INNER JOIN `rocpd_kernel_dispatch` K ON K.event_id = PMC_E.event_id + AND K.guid = PMC_E.guid + INNER JOIN `rocpd_info_process` P ON P.id = K.pid + AND P.guid = K.guid; + +-- events with arguments --- +CREATE VIEW IF NOT EXISTS + `events_args` AS +SELECT + E.id AS event_id, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + E.stack_id, + E.parent_stack_id, + E.correlation_id, + A.position AS arg_position, + A.type AS arg_type, + A.name AS arg_name, + A.value AS arg_value, + E.call_stack, + E.line_info, + A.extdata +FROM + `rocpd_event` E + INNER JOIN `rocpd_arg` A ON A.event_id = E.id + AND A.guid = E.guid; + +-- list of astream arguments enriched by the corresponding stream descriptions +CREATE VIEW IF NOT EXISTS + `stream_args` AS +SELECT + A.id AS argument_id, + A.event_id AS event_id, + A.position AS arg_position, + A.type AS arg_type, + A.value AS arg_value, + JSON_EXTRACT(A.extdata, '$.stream_id') AS stream_id, + S.nid, + P.pid, + S.name AS stream_name, + S.extdata AS extdata +FROM + `rocpd_arg` A + INNER JOIN `rocpd_info_stream` S ON JSON_EXTRACT(A.extdata, '$.stream_id') = S.id + AND A.guid = S.guid + INNER JOIN `rocpd_info_process` P ON P.id = S.pid + AND P.guid = S.guid +WHERE + A.name = 'stream'; + +-- +-- +CREATE VIEW IF NOT EXISTS + `memory_copies` AS +SELECT + M.id, + M.guid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + M.nid, + P.pid, + T.tid, + M.start, + M.end, + (M.end - M.start) AS duration, + S.string AS name, + R.string AS region_name, + M.stream_id, + M.queue_id, + M.graph_exec_id, + M.graph_node_id, + ST.name AS stream_name, + Q.name AS queue_name, + M.size, + dst_agent.name AS dst_device, + dst_agent.absolute_index AS dst_agent_abs_index, + dst_agent.logical_index AS dst_agent_log_index, + dst_agent.type_index AS dst_agent_type_index, + dst_agent.type AS dst_agent_type, + M.dst_address, + src_agent.name AS src_device, + src_agent.absolute_index AS src_agent_abs_index, + src_agent.logical_index AS src_agent_log_index, + src_agent.type_index AS src_agent_type_index, + src_agent.type AS src_agent_type, + M.src_address, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id +FROM + `rocpd_memory_copy` M + INNER JOIN `rocpd_string` S ON S.id = M.name_id + AND S.guid = M.guid + LEFT JOIN `rocpd_string` R ON R.id = M.region_name_id + AND R.guid = M.guid + INNER JOIN `rocpd_info_agent` dst_agent ON dst_agent.id = M.dst_agent_id + AND dst_agent.guid = M.guid + INNER JOIN `rocpd_info_agent` src_agent ON src_agent.id = M.src_agent_id + AND src_agent.guid = M.guid + LEFT JOIN `rocpd_info_queue` Q ON Q.id = M.queue_id + AND Q.guid = M.guid + LEFT JOIN `rocpd_info_stream` ST ON ST.id = M.stream_id + AND ST.guid = M.guid + INNER JOIN `rocpd_event` E ON E.id = M.event_id + AND E.guid = M.guid + INNER JOIN `rocpd_info_process` P ON P.id = M.pid + AND P.guid = M.guid + INNER JOIN `rocpd_info_thread` T ON T.id = M.tid + AND T.guid = M.guid; + +-- +-- HIP event barrier records +CREATE VIEW IF NOT EXISTS + `hip_events` AS +SELECT + H.id, + H.guid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + H.nid, + P.pid, + T.tid, + H.start, + H.end, + (H.end - H.start) AS duration, + S.string AS name, + H.stream_id, + H.queue_id, + ST.name AS stream_name, + Q.name AS queue_name, + agent.name AS device, + agent.absolute_index AS agent_abs_index, + agent.logical_index AS agent_log_index, + agent.type_index AS agent_type_index, + agent.type AS agent_type, + H.hip_event_handle, + H.source_queue_id, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id +FROM + `rocpd_hip_event` H + INNER JOIN `rocpd_string` S ON S.id = H.name_id + AND S.guid = H.guid + LEFT JOIN `rocpd_info_agent` agent ON agent.id = H.agent_id + AND agent.guid = H.guid + LEFT JOIN `rocpd_info_queue` Q ON Q.id = H.queue_id + AND Q.guid = H.guid + LEFT JOIN `rocpd_info_stream` ST ON ST.id = H.stream_id + AND ST.guid = H.guid + INNER JOIN `rocpd_event` E ON E.id = H.event_id + AND E.guid = H.guid + INNER JOIN `rocpd_info_process` P ON P.id = H.pid + AND P.guid = H.guid + INNER JOIN `rocpd_info_thread` T ON T.id = H.tid + AND T.guid = H.guid; + +-- +-- HIP graph launch summary records +CREATE VIEW IF NOT EXISTS + `graph_launches` AS +SELECT + G.id, + G.guid, + G.nid, + P.pid, + T.tid, + G.agent_id, + A.absolute_index AS agent_abs_index, + A.logical_index AS agent_log_index, + A.type_index AS agent_type_index, + A.type AS agent_type, + G.queue_id, + Q.name AS queue, + 'Graph Execution' AS name, + G.start, + G.end, + (G.end - G.start) AS duration, + G.graph_exec_id, + G.kernel_dispatch_count, + G.event_id, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id +FROM + `rocpd_graph_launch` G + LEFT JOIN `rocpd_info_agent` A ON A.id = G.agent_id + AND A.guid = G.guid + LEFT JOIN `rocpd_info_queue` Q ON Q.id = G.queue_id + AND Q.guid = G.guid + INNER JOIN `rocpd_event` E ON E.id = G.event_id + AND E.guid = G.guid + INNER JOIN `rocpd_info_process` P ON P.id = G.pid + AND P.guid = G.guid + INNER JOIN `rocpd_info_thread` T ON T.id = G.tid + AND T.guid = G.guid; + +-- +-- +CREATE VIEW IF NOT EXISTS + `memory_allocations` AS +SELECT + M.id, + M.guid, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + M.nid, + P.pid, + T.tid, + M.start, + M.end, + (M.end - M.start) AS duration, + M.type, + M.level, + A.name AS agent_name, + A.absolute_index AS agent_abs_index, + A.logical_index AS agent_log_index, + A.type_index AS agent_type_index, + A.type AS agent_type, + M.address, + M.size, + M.queue_id, + Q.name AS queue_name, + M.stream_id, + ST.name AS stream_name, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id +FROM + `rocpd_memory_allocate` M + LEFT JOIN `rocpd_info_agent` A ON M.agent_id = A.id + AND M.guid = A.guid + LEFT JOIN `rocpd_info_queue` Q ON Q.id = M.queue_id + AND Q.guid = M.guid + LEFT JOIN `rocpd_info_stream` ST ON ST.id = M.stream_id + AND ST.guid = M.guid + INNER JOIN `rocpd_event` E ON E.id = M.event_id + AND E.guid = M.guid + INNER JOIN `rocpd_info_process` P ON P.id = M.pid + AND P.guid = M.guid + INNER JOIN `rocpd_info_thread` T ON T.id = M.tid + AND P.guid = M.guid; + +-- +-- +CREATE VIEW IF NOT EXISTS + `scratch_memory` AS +SELECT + M.id, + M.guid, + M.nid, + P.pid, + M.type AS operation, + A.name AS agent_name, + A.absolute_index AS agent_abs_index, + A.logical_index AS agent_log_index, + A.type_index AS agent_type_index, + A.type AS agent_type, + M.queue_id, + T.tid, + JSON_EXTRACT(M.extdata, '$.flags') AS alloc_flags, + M.start, + M.end, + (M.end - M.start) AS duration, + M.size, + M.address, + E.correlation_id, + E.stack_id, + E.parent_stack_id, + E.correlation_id AS corr_id, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + E.extdata AS event_extdata +FROM + `rocpd_memory_allocate` M + LEFT JOIN `rocpd_info_agent` A ON M.agent_id = A.id + AND M.guid = A.guid + LEFT JOIN `rocpd_info_queue` Q ON Q.id = M.queue_id + AND Q.guid = M.guid + INNER JOIN `rocpd_event` E ON E.id = M.event_id + AND E.guid = M.guid + INNER JOIN `rocpd_info_process` P ON P.id = M.pid + AND P.guid = M.guid + INNER JOIN `rocpd_info_thread` T ON T.id = M.tid + AND T.guid = M.guid +WHERE + M.level = 'SCRATCH' +ORDER BY + M.start ASC; + +-- +-- +CREATE VIEW IF NOT EXISTS + `counters_collection` AS +SELECT + MIN(PMC_E.id) AS id, + PMC_E.guid, + K.dispatch_id, + K.kernel_id, + E.id AS event_id, + E.correlation_id, + E.stack_id, + E.parent_stack_id, + P.pid, + T.tid, + K.agent_id, + A.absolute_index AS agent_abs_index, + A.logical_index AS agent_log_index, + A.type_index AS agent_type_index, + A.type AS agent_type, + K.queue_id, + K.grid_size_x AS grid_size_x, + K.grid_size_y AS grid_size_y, + K.grid_size_z AS grid_size_z, + (K.grid_size_x * K.grid_size_y * K.grid_size_z) AS grid_size, + KS.display_name AS kernel_name, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = K.region_name_id + AND RS.guid = K.guid + ) AS kernel_region, + K.workgroup_size_x AS workgroup_size_x, + K.workgroup_size_y AS workgroup_size_y, + K.workgroup_size_z AS workgroup_size_z, + (K.workgroup_size_x * K.workgroup_size_y * K.workgroup_size_z) AS workgroup_size, + K.group_segment_size AS lds_block_size, + K.private_segment_size AS scratch_size, + KS.arch_vgpr_count AS vgpr_count, + KS.accum_vgpr_count, + KS.sgpr_count, + PMC_I.name AS counter_name, + PMC_I.symbol AS counter_symbol, + PMC_I.component, + PMC_I.description, + PMC_I.block, + PMC_I.expression, + PMC_I.value_type, + PMC_I.id AS counter_id, + SUM(PMC_E.value) AS value, + K.start, + K.end, + PMC_I.is_constant, + PMC_I.is_derived, + (K.end - K.start) AS duration, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + K.nid, + E.extdata, + KS.code_object_id +FROM + `rocpd_pmc_event` PMC_E + INNER JOIN `rocpd_info_pmc` PMC_I ON PMC_I.id = PMC_E.pmc_id + AND PMC_I.guid = PMC_E.guid + INNER JOIN `rocpd_event` E ON E.id = PMC_E.event_id + AND E.guid = PMC_E.guid + INNER JOIN `rocpd_kernel_dispatch` K ON K.event_id = PMC_E.event_id + AND K.guid = PMC_E.guid + INNER JOIN `rocpd_info_agent` A ON A.id = K.agent_id + AND A.guid = K.guid + INNER JOIN `rocpd_info_kernel_symbol` KS ON KS.id = K.kernel_id + AND KS.guid = K.guid + INNER JOIN `rocpd_info_process` P ON P.id = K.pid + AND P.guid = K.guid + INNER JOIN `rocpd_info_thread` T ON T.id = K.tid + AND T.guid = K.guid +WHERE + PMC_E.sample_id IS NULL +GROUP BY + PMC_E.guid, + K.dispatch_id, + PMC_I.name, + K.agent_id; + +-- +-- SPM counters view: no aggregation, preserves per-SE records +-- +CREATE VIEW IF NOT EXISTS + `spm_counters` AS +SELECT + PMC_E.id, + PMC_E.guid, + K.dispatch_id, + K.kernel_id, + E.id AS event_id, + E.correlation_id, + E.stack_id, + E.parent_stack_id, + P.pid, + T.tid, + K.agent_id, + A.absolute_index AS agent_abs_index, + A.logical_index AS agent_log_index, + A.type_index AS agent_type_index, + A.type AS agent_type, + K.queue_id, + K.grid_size_x AS grid_size_x, + K.grid_size_y AS grid_size_y, + K.grid_size_z AS grid_size_z, + (K.grid_size_x * K.grid_size_y * K.grid_size_z) AS grid_size, + KS.display_name AS kernel_name, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = K.region_name_id + AND RS.guid = K.guid + ) AS kernel_region, + K.workgroup_size_x AS workgroup_size_x, + K.workgroup_size_y AS workgroup_size_y, + K.workgroup_size_z AS workgroup_size_z, + (K.workgroup_size_x * K.workgroup_size_y * K.workgroup_size_z) AS workgroup_size, + K.group_segment_size AS lds_block_size, + K.private_segment_size AS scratch_size, + KS.arch_vgpr_count AS vgpr_count, + KS.accum_vgpr_count, + KS.sgpr_count, + PMC_I.name AS counter_name, + PMC_I.symbol AS counter_symbol, + PMC_I.component, + PMC_I.description, + PMC_I.block, + PMC_I.expression, + PMC_I.value_type, + PMC_I.id AS counter_id, + PMC_I.xcc, + PMC_I.shader_engine, + PMC_I.shader_array, + PMC_I.wgp, + PMC_I.instance, + PMC_E.value, + S.timestamp, + PMC_I.is_constant, + PMC_I.is_derived, + ( + SELECT + string + FROM + `rocpd_string` RS + WHERE + RS.id = E.category_id + AND RS.guid = E.guid + ) AS category, + K.nid, + E.extdata, + KS.code_object_id +FROM + `rocpd_pmc_event` PMC_E + INNER JOIN `rocpd_sample` S ON S.id = PMC_E.sample_id + AND S.guid = PMC_E.guid + INNER JOIN `rocpd_info_pmc` PMC_I ON PMC_I.id = PMC_E.pmc_id + AND PMC_I.guid = PMC_E.guid + INNER JOIN `rocpd_event` E ON E.id = PMC_E.event_id + AND E.guid = PMC_E.guid + INNER JOIN `rocpd_kernel_dispatch` K ON K.event_id = PMC_E.event_id + AND K.guid = PMC_E.guid + INNER JOIN `rocpd_info_agent` A ON A.id = K.agent_id + AND A.guid = K.guid + INNER JOIN `rocpd_info_kernel_symbol` KS ON KS.id = K.kernel_id + AND KS.guid = K.guid + INNER JOIN `rocpd_info_process` P ON P.id = K.pid + AND P.guid = K.guid + INNER JOIN `rocpd_info_thread` T ON T.id = K.tid + AND T.guid = K.guid; diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_indexes.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_indexes.sql new file mode 100644 index 000000000000..43d722ea10e4 --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_indexes.sql @@ -0,0 +1,45 @@ +-- +-- Indexes for the various fields +-- + +-- string field +-- CREATE INDEX `rocpd_string{{uuid}}_string_idx` ON `rocpd_string{{uuid}}` ("string"); + +-- guid field +-- CREATE INDEX `rocpd_string{{uuid}}_guid_idx` ON `rocpd_string{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_node{{uuid}}_guid_idx` ON `rocpd_info_node{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_process{{uuid}}_guid_idx` ON `rocpd_info_process{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_thread{{uuid}}_guid_idx` ON `rocpd_info_thread{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_agent{{uuid}}_guid_idx` ON `rocpd_info_agent{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_queue{{uuid}}_guid_idx` ON `rocpd_info_queue{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_stream{{uuid}}_guid_idx` ON `rocpd_info_stream{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_pmc{{uuid}}_guid_idx` ON `rocpd_info_pmc{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_code_object{{uuid}}_guid_idx` ON `rocpd_info_code_object{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_info_kernel_symbol{{uuid}}_guid_idx` ON `rocpd_info_kernel_symbol{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_track{{uuid}}_guid_idx` ON `rocpd_track{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_event{{uuid}}_guid_idx` ON `rocpd_event{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_arg{{uuid}}_guid_idx` ON `rocpd_arg{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_pmc_event{{uuid}}_guid_idx` ON `rocpd_pmc_event{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_region{{uuid}}_guid_idx` ON `rocpd_region{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_sample{{uuid}}_guid_idx` ON `rocpd_sample{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_kernel_dispatch{{uuid}}_guid_idx` ON `rocpd_kernel_dispatch{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_memory_copy{{uuid}}_guid_idx` ON `rocpd_memory_copy{{uuid}}` ("id", "guid"); +-- CREATE INDEX `rocpd_memory_allocate{{uuid}}_guid_idx` ON `rocpd_memory_allocate{{uuid}}` ("id", "guid"); + +-- CREATE INDEX `rocpd_event{{uuid}}_category_idx` ON `rocpd_event{{uuid}}` ("id", "guid", "category_id"); +-- CREATE INDEX `rocpd_region{{uuid}}_event_idx` ON `rocpd_region{{uuid}}` ("id", "guid", "event_id"); +-- CREATE INDEX `rocpd_region{{uuid}}_name_idx` ON `rocpd_region{{uuid}}` ("id", "guid", "name_id"); +-- CREATE INDEX `rocpd_sample{{uuid}}_event_idx` ON `rocpd_sample{{uuid}}` ("id", "guid", "event_id"); +-- CREATE INDEX `rocpd_sample{{uuid}}_track_idx` ON `rocpd_sample{{uuid}}` ("id", "guid", "track_id"); +-- CREATE INDEX `rocpd_track{{uuid}}_name_idx` ON `rocpd_track{{uuid}}` ("id", "guid", "name_id"); + +-- CREATE INDEX `rocpd_memory_copy{{uuid}}_guid_nid_pid_idx` ON `rocpd_memory_copy{{uuid}}` ("guid", "nid", "pid"); +-- CREATE INDEX `rocpd_kernel_dispatch{{uuid}}_guid_nid_pid_idx` ON `rocpd_kernel_dispatch{{uuid}}` ("guid", "nid", "pid"); +-- CREATE INDEX `rocpd_region{{uuid}}_guid_idx` ON `rocpd_region{{uuid}}` ("guid", "nid", "pid"); +-- CREATE INDEX `rocpd_sample{{uuid}}_guid_nid_pid_idx` ON `rocpd_sample{{uuid}}` ("guid", "nid", "pid"); + +-- CREATE INDEX `rocpd_region{{uuid}}_guid_idx` ON `rocpd_region{{uuid}}` ("guid"); +-- CREATE INDEX `rocpd_region{{uuid}}_nid_idx` ON `rocpd_region{{uuid}}` ("nid"); +-- CREATE INDEX `rocpd_region{{uuid}}_pid_idx` ON `rocpd_region{{uuid}}` ("pid"); +-- CREATE INDEX `rocpd_region{{uuid}}_start_idx` ON `rocpd_region{{uuid}}` ("start"); +-- CREATE INDEX `rocpd_region{{uuid}}_end_idx` ON `rocpd_region{{uuid}}` ("end"); diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_metadata.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_metadata.sql new file mode 100644 index 000000000000..904bca4f606a --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_metadata.sql @@ -0,0 +1,14 @@ +-- +-- Standard metadata insertion +-- +-- +INSERT INTO + `rocpd_metadata{{uuid}}` ("tag", "value") +VALUES + ("schema_version", "{{schema_version}}"), -- full version string for the current version + ("schema_version_major", "{{schema_version_major}}"), -- major ID for the current version + ("schema_version_minor", "{{schema_version_minor}}"), -- minor ID for the current version + ("schema_version_patch", "{{schema_version_patch}}"), -- patch ID for the current version + ("schema_creation_time", CURRENT_TIMESTAMP), -- time of schema creation + ("uuid", "{{uuid}}"), + ("guid", "{{guid}}"); diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_tables.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_tables.sql new file mode 100644 index 000000000000..4133c12455c7 --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_tables.sql @@ -0,0 +1,430 @@ +-- RocPD schema version 3.0.5 + +CREATE TABLE IF NOT EXISTS + "rocpd_metadata{{uuid}}" ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "tag" TEXT NOT NULL, + "value" TEXT NOT NULL + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_string{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "string" TEXT NOT NULL UNIQUE ON CONFLICT ABORT + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_node{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "hash" BIGINT NOT NULL UNIQUE, + "machine_id" TEXT NOT NULL UNIQUE, + "system_name" TEXT, + "hostname" TEXT, + "release" TEXT, + "version" TEXT, + "hardware_name" TEXT, + "domain_name" TEXT + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_process{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "ppid" INTEGER, + "pid" INTEGER NOT NULL, + "init" BIGINT, + "fini" BIGINT, + "start" BIGINT, + "end" BIGINT, + "command" TEXT, + "environment" JSONB DEFAULT "{}" NOT NULL, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_thread{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "ppid" INTEGER, + "pid" INTEGER NOT NULL, + "tid" INTEGER NOT NULL, + "name" TEXT, + "start" BIGINT, + "end" BIGINT, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_agent{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "type" TEXT CHECK ("type" IN ('CPU', 'GPU', 'NIC')), + "absolute_index" INTEGER, + "logical_index" INTEGER, + "type_index" INTEGER, + "uuid" INTEGER, + "name" TEXT, + "model_name" TEXT, + "vendor_name" TEXT, + "product_name" TEXT, + "user_name" TEXT, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_queue{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "name" TEXT, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_stream{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "name" TEXT, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Performance monitoring counters (PMC) descriptions +CREATE TABLE IF NOT EXISTS + `rocpd_info_pmc{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "agent_id" INTEGER, + "target_arch" TEXT CHECK ("target_arch" IN ('CPU', 'GPU', 'NIC')), + "event_code" INT, + "instance_id" INTEGER, + "name" TEXT NOT NULL, + "symbol" TEXT NOT NULL, + "description" TEXT, + "long_description" TEXT DEFAULT "", + "component" TEXT, + "units" TEXT DEFAULT "", + "value_type" TEXT CHECK ("value_type" IN ('ABS', 'ACCUM', 'RELATIVE')), + "block" TEXT, + "expression" TEXT, + "is_constant" INTEGER, + "is_derived" INTEGER, + "spm_support" INTEGER, + -- SPM location. NULL on the base counter row. On a location row, xcc, + -- shader_engine, shader_array, and wgp are NULL when the counter does + -- not provide that dimension. instance is always set on a location row. + "xcc" INTEGER, + "shader_engine" INTEGER, + "shader_array" INTEGER, + "wgp" INTEGER, + "instance" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_code_object{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "agent_id" INTEGER, + "uri" TEXT, + "load_base" BIGINT, + "load_size" BIGINT, + "load_delta" BIGINT, + "storage_type" TEXT CHECK ("storage_type" IN ('FILE', 'MEMORY')), + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_info_kernel_symbol{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "code_object_id" INTEGER NOT NULL, + "kernel_name" TEXT, + "display_name" TEXT, + "kernel_object" INTEGER, + "kernarg_segment_size" INTEGER, + "kernarg_segment_alignment" INTEGER, + "group_segment_size" INTEGER, + "private_segment_size" INTEGER, + "sgpr_count" INTEGER, + "arch_vgpr_count" INTEGER, + "accum_vgpr_count" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (code_object_id) REFERENCES `rocpd_info_code_object{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Stores repetitive info for samples +CREATE TABLE IF NOT EXISTS + `rocpd_track{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER, + "tid" INTEGER, + "name_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (name_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Storage for a region, instant, and counter +CREATE TABLE IF NOT EXISTS + `rocpd_event{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "category_id" INTEGER, + "stack_id" INTEGER, + "parent_stack_id" INTEGER, + "correlation_id" INTEGER, + "call_stack" JSONB DEFAULT "{}" NOT NULL, + "line_info" JSONB DEFAULT "{}" NOT NULL, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (category_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- stores arguments for events +CREATE TABLE IF NOT EXISTS + `rocpd_arg{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "event_id" INTEGER NOT NULL, + "position" INTEGER NOT NULL, + "type" TEXT NOT NULL, + "name" TEXT NOT NULL, + "value" TEXT, -- TODO: discuss make it value_id and integer, refer to string table -- + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Region with a start/stop on the same thread (CPU) +CREATE TABLE IF NOT EXISTS + `rocpd_pmc_event{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "event_id" INTEGER, + "sample_id" INTEGER, -- NULL for PMC, non-NULL for SPM (timestamped streaming counters) + "pmc_id" INTEGER NOT NULL, + "value" REAL DEFAULT 0.0, + "extdata" JSONB DEFAULT "{}", + FOREIGN KEY (pmc_id) REFERENCES `rocpd_info_pmc{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (sample_id) REFERENCES `rocpd_sample{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Region with a start/stop on the same thread (CPU) +CREATE TABLE IF NOT EXISTS + `rocpd_region{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "tid" INTEGER NOT NULL, + "start" BIGINT NOT NULL, + "end" BIGINT NOT NULL, + "name_id" INTEGER NOT NULL, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (name_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Instantaneous sample +CREATE TABLE IF NOT EXISTS + `rocpd_sample{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "track_id" INTEGER NOT NULL, + "timestamp" BIGINT NOT NULL, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (track_id) REFERENCES `rocpd_track{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_kernel_dispatch{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "tid" INTEGER, + "agent_id" INTEGER NOT NULL, + "kernel_id" INTEGER NOT NULL, + "dispatch_id" INTEGER NOT NULL, + "queue_id" INTEGER NOT NULL, + "stream_id" INTEGER NOT NULL, + "start" BIGINT NOT NULL, + "end" BIGINT NOT NULL, + "private_segment_size" INTEGER, + "group_segment_size" INTEGER, + "workgroup_size_x" INTEGER NOT NULL, + "workgroup_size_y" INTEGER NOT NULL, + "workgroup_size_z" INTEGER NOT NULL, + "grid_size_x" INTEGER NOT NULL, + "grid_size_y" INTEGER NOT NULL, + "grid_size_z" INTEGER NOT NULL, + "graph_exec_id" INTEGER NOT NULL DEFAULT 0, + "graph_node_id" INTEGER NOT NULL DEFAULT 0, + "region_name_id" INTEGER, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (kernel_id) REFERENCES `rocpd_info_kernel_symbol{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (queue_id) REFERENCES `rocpd_info_queue{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (stream_id) REFERENCES `rocpd_info_stream{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (region_name_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +CREATE TABLE IF NOT EXISTS + `rocpd_memory_copy{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "tid" INTEGER, + "start" BIGINT NOT NULL, + "end" BIGINT NOT NULL, + "name_id" INTEGER NOT NULL, + "dst_agent_id" INTEGER, + "dst_address" INTEGER, + "src_agent_id" INTEGER, + "src_address" INTEGER, + "size" INTEGER NOT NULL, + "queue_id" INTEGER, + "stream_id" INTEGER, + "graph_exec_id" INTEGER NOT NULL DEFAULT 0, + "graph_node_id" INTEGER NOT NULL DEFAULT 0, + "region_name_id" INTEGER, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (name_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (dst_agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (src_agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (stream_id) REFERENCES `rocpd_info_stream{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (queue_id) REFERENCES `rocpd_info_queue{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (region_name_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- Memory allocations (real memory, virtual memory, and scratch memory) +CREATE TABLE IF NOT EXISTS + `rocpd_memory_allocate{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "tid" INTEGER, + "agent_id" INTEGER, + "type" TEXT CHECK ("type" IN ('ALLOC', 'FREE', 'REALLOC', 'RECLAIM')), + "level" TEXT CHECK ("level" IN ('REAL', 'VIRTUAL', 'SCRATCH')), + "start" BIGINT NOT NULL, + "end" BIGINT NOT NULL, + "address" INTEGER, + "size" INTEGER NOT NULL, + "queue_id" INTEGER, + "stream_id" INTEGER, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (stream_id) REFERENCES `rocpd_info_stream{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (queue_id) REFERENCES `rocpd_info_queue{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- HIP event barrier records (hipEventRecord and hipStreamWaitEvent completions). +CREATE TABLE IF NOT EXISTS + `rocpd_hip_event{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "tid" INTEGER, + "start" BIGINT NOT NULL, + "end" BIGINT NOT NULL, + "name_id" INTEGER NOT NULL, + "agent_id" INTEGER, + "queue_id" INTEGER, + "stream_id" INTEGER, + "hip_event_handle" BIGINT NOT NULL, + "source_queue_id" INTEGER, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (name_id) REFERENCES `rocpd_string{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (stream_id) REFERENCES `rocpd_info_stream{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (queue_id) REFERENCES `rocpd_info_queue{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); + +-- HIP graph launch summary records (one per successful hipGraphLaunch). +CREATE TABLE IF NOT EXISTS + `rocpd_graph_launch{{uuid}}` ( + "id" INTEGER NOT NULL PRIMARY KEY AUTOINCREMENT, + "guid" TEXT DEFAULT "{{guid}}" NOT NULL, + "nid" INTEGER NOT NULL, + "pid" INTEGER NOT NULL, + "tid" INTEGER, + "agent_id" INTEGER, + "queue_id" INTEGER, + "start" BIGINT NOT NULL, + "end" BIGINT NOT NULL, + "graph_exec_id" INTEGER NOT NULL, + "kernel_dispatch_count" INTEGER NOT NULL DEFAULT 0, + "event_id" INTEGER, + "extdata" JSONB DEFAULT "{}" NOT NULL, + FOREIGN KEY (nid) REFERENCES `rocpd_info_node{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (pid) REFERENCES `rocpd_info_process{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (tid) REFERENCES `rocpd_info_thread{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (agent_id) REFERENCES `rocpd_info_agent{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (queue_id) REFERENCES `rocpd_info_queue{{uuid}}` (id) ON UPDATE CASCADE, + FOREIGN KEY (event_id) REFERENCES `rocpd_event{{uuid}}` (id) ON UPDATE CASCADE + ); diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_views.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_views.sql new file mode 100644 index 000000000000..5925251d1fe1 --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/rocpd_views.sql @@ -0,0 +1,153 @@ +CREATE VIEW IF NOT EXISTS + `rocpd_metadata` AS +SELECT + * +FROM + `rocpd_metadata{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_string` AS +SELECT + * +FROM + `rocpd_string{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_node` AS +SELECT + * +FROM + `rocpd_info_node{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_process` AS +SELECT + * +FROM + `rocpd_info_process{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_thread` AS +SELECT + * +FROM + `rocpd_info_thread{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_agent` AS +SELECT + * +FROM + `rocpd_info_agent{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_queue` AS +SELECT + * +FROM + `rocpd_info_queue{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_stream` AS +SELECT + * +FROM + `rocpd_info_stream{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_pmc` AS +SELECT + * +FROM + `rocpd_info_pmc{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_code_object` AS +SELECT + * +FROM + `rocpd_info_code_object{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_info_kernel_symbol` AS +SELECT + * +FROM + `rocpd_info_kernel_symbol{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_track` AS +SELECT + * +FROM + `rocpd_track{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_event` AS +SELECT + * +FROM + `rocpd_event{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_arg` AS +SELECT + * +FROM + `rocpd_arg{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_pmc_event` AS +SELECT + * +FROM + `rocpd_pmc_event{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_region` AS +SELECT + * +FROM + `rocpd_region{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_sample` AS +SELECT + * +FROM + `rocpd_sample{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_kernel_dispatch` AS +SELECT + * +FROM + `rocpd_kernel_dispatch{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_memory_copy` AS +SELECT + * +FROM + `rocpd_memory_copy{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_memory_allocate` AS +SELECT + * +FROM + `rocpd_memory_allocate{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_hip_event` AS +SELECT + * +FROM + `rocpd_hip_event{{uuid}}`; + +CREATE VIEW IF NOT EXISTS + `rocpd_graph_launch` AS +SELECT + * +FROM + `rocpd_graph_launch{{uuid}}`; diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/summary_views.sql b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/summary_views.sql new file mode 100644 index 000000000000..b4c713fe3a84 --- /dev/null +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/3.0.5/summary_views.sql @@ -0,0 +1,153 @@ +-- +-- Useful summary views +-- +-- +-- Sorted list of kernels which consume the most overall time +CREATE VIEW IF NOT EXISTS + `top_kernels` AS +SELECT + S.display_name AS name, + COUNT(K.kernel_id) AS total_calls, + SUM(K.end - K.start) / 1000.0 AS total_duration, + (SUM(K.end - K.start) / COUNT(K.kernel_id)) / 1000.0 AS average, + SUM(K.end - K.start) * 100.0 / NULLIF(( + SELECT + SUM(A.end - A.start) + FROM + `rocpd_kernel_dispatch` A + ), 0) AS percentage +FROM + `rocpd_kernel_dispatch` K + INNER JOIN `rocpd_info_kernel_symbol` S ON S.id = K.kernel_id + AND S.guid = K.guid +GROUP BY + name +ORDER BY + total_duration DESC; + +-- +-- GPU utilization metrics including kernels and memory copy operations +CREATE VIEW IF NOT EXISTS + `busy` AS +SELECT + A.agent_id, + AG.type, + GpuTime, + WallTime, + GpuTime * 1.0 / NULLIF(WallTime, 0) AS Busy +FROM + ( + SELECT + agent_id, + guid, + SUM(END - start) AS GpuTime + FROM + ( + SELECT + agent_id, + guid, + END, + start + FROM + `rocpd_kernel_dispatch` + UNION ALL + SELECT + dst_agent_id AS agent_id, + guid, + END, + start + FROM + `rocpd_memory_copy` + ) + GROUP BY + agent_id, + guid + ) A + INNER JOIN ( + SELECT + MAX(END) - MIN(start) AS WallTime + FROM + ( + SELECT + END, + start + FROM + `rocpd_kernel_dispatch` + UNION ALL + SELECT + END, + start + FROM + `rocpd_memory_copy` + ) + ) W ON 1 = 1 + INNER JOIN `rocpd_info_agent` AG ON AG.id = A.agent_id + AND AG.guid = A.guid; + +-- +-- Overall performance summary including kernels and memory copy operations +CREATE VIEW + `top` AS +SELECT + name, + COUNT(*) AS total_calls, + SUM(duration) / 1000.0 AS total_duration, + (SUM(duration) / COUNT(*)) / 1000.0 AS average, + SUM(duration) * 100.0 / NULLIF(total_time, 0) AS percentage +FROM + ( + -- Kernel operations + SELECT + ks.display_name AS name, + (kd.end - kd.start) AS duration + FROM + `rocpd_kernel_dispatch` kd + INNER JOIN `rocpd_info_kernel_symbol` ks ON kd.kernel_id = ks.id + AND kd.guid = ks.guid + UNION ALL + -- Memory operations + SELECT + rs.string AS name, + (END - start) AS duration + FROM + `rocpd_memory_copy` mc + INNER JOIN `rocpd_string` rs ON rs.id = mc.name_id + AND rs.guid = mc.guid + UNION ALL + -- Regions + SELECT + rs.string AS name, + (END - start) AS duration + FROM + `rocpd_region` rr + INNER JOIN `rocpd_string` rs ON rs.id = rr.name_id + AND rs.guid = rr.guid + ) operations + CROSS JOIN ( + SELECT + SUM(END - start) AS total_time + FROM + ( + SELECT + END, + start + FROM + `rocpd_kernel_dispatch` + UNION ALL + SELECT + END, + start + FROM + `rocpd_memory_copy` + UNION ALL + SELECT + END, + start + FROM + `rocpd_region` + ) + ) TOTAL +GROUP BY + name +ORDER BY + total_duration DESC; diff --git a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/CMakeLists.txt b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/CMakeLists.txt index e8fd7ac83d4f..6dfda8bd3efe 100644 --- a/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/CMakeLists.txt +++ b/projects/rocprofiler-sdk/source/share/rocprofiler-sdk-rocpd/versions/CMakeLists.txt @@ -7,3 +7,4 @@ add_subdirectory(3.0.1) add_subdirectory(3.0.2) add_subdirectory(3.0.3) add_subdirectory(3.0.4) +add_subdirectory(3.0.5) diff --git a/projects/rocprofiler-sdk/tests/rocpd-back-compat/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocpd-back-compat/CMakeLists.txt index 4f17857b2a0e..a7f912fb2213 100644 --- a/projects/rocprofiler-sdk/tests/rocpd-back-compat/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocpd-back-compat/CMakeLists.txt @@ -71,7 +71,7 @@ find_package(Python3 REQUIRED COMPONENTS Interpreter) # ships, append it here and add the corresponding validate assertions in validate.py. Add # appropriate changes to make_db.py to insert dummy data into a fake DB to simulate the # new schema changes. -set(ROCPD_COMPAT_TESTED_SCHEMAS "3.0.0" "3.0.1" "3.0.2" "3.0.3" "3.0.4" "3.0.5") +set(ROCPD_COMPAT_TESTED_SCHEMAS "3.0.0" "3.0.1" "3.0.2" "3.0.3" "3.0.4" "3.0.5" "3.0.6") # Derive old schemas (all except the latest) and the latest schema. list(LENGTH ROCPD_COMPAT_TESTED_SCHEMAS _num_schemas) @@ -102,12 +102,24 @@ endif() # --------------------------------------------------------------------------- set(rocpd-back-compat-schema-path "${CMAKE_BINARY_DIR}/share/rocprofiler-sdk-rocpd") +# Prefer the configure-time Python's site-packages, but also put every +# python*/site-packages under the build lib dir on PYTHONPATH. A mismatch between +# Python3_EXECUTABLE and the directory that received the rocpd package (seen when the +# build tree has python3.10 packages but tests were configured against python3.12) +# otherwise makes make_db import a system rocpd and write the wrong schema views into the +# fixture databases. set(rocpd-back-compat-site-packages "${rocprofiler-sdk_LIB_DIR}/python${Python3_VERSION_MAJOR}.${Python3_VERSION_MINOR}/site-packages" ) +file(GLOB _rocpd_back_compat_all_site_packages + "${rocprofiler-sdk_LIB_DIR}/python*/site-packages") +list(INSERT _rocpd_back_compat_all_site_packages 0 "${rocpd-back-compat-site-packages}") +list(REMOVE_DUPLICATES _rocpd_back_compat_all_site_packages) +string(REPLACE ";" ":" rocpd-back-compat-site-packages-path + "${_rocpd_back_compat_all_site_packages}") set(rocpd-back-compat-env "ROCPD_SCHEMA_PATH=${rocpd-back-compat-schema-path}" - "PYTHONPATH=${ROCPROFILER_SDK_TESTS_BINARY_DIR}/pytest-packages:${rocpd-back-compat-site-packages}" + "PYTHONPATH=${ROCPROFILER_SDK_TESTS_BINARY_DIR}/pytest-packages:${rocpd-back-compat-site-packages-path}" ) # Root directory under which per-version output subdirectories are created: @@ -211,7 +223,7 @@ rocprofiler_add_integration_validate_test( LABELS "integration-tests;rocpd" PRELOAD "${ROCPROFILER_MEMCHECK_PRELOAD_ENV_VALUE}" ENVIRONMENT "ROCPD_SCHEMA_PATH=${rocpd-back-compat-schema-path}" - PYTHON_PATH "${rocpd-back-compat-site-packages}" + PYTHON_PATH "${rocpd-back-compat-site-packages-path}" FIXTURES_REQUIRED ${_gen_fixtures} DEPENDS rocprofiler-sdk::rocprofv3 ARGS --output-root ${ROCPD_BACK_COMPAT_OUTPUT_DIR} --tested-schemas diff --git a/projects/rocprofiler-sdk/tests/rocpd-back-compat/make_db.py b/projects/rocprofiler-sdk/tests/rocpd-back-compat/make_db.py index e802c6da378d..06310b7d3cd8 100644 --- a/projects/rocprofiler-sdk/tests/rocpd-back-compat/make_db.py +++ b/projects/rocprofiler-sdk/tests/rocpd-back-compat/make_db.py @@ -74,6 +74,8 @@ _PMC_NO_SPM_ID = 1 _PMC_SPM_ID = 2 +_PMC_GRBM_COUNT_ID = 3 +_PMC_SQ_INSTS_SALU_ID = 4 _TRACK_SPM_ID = 1 _SAMPLE_SPM_ID = 1 @@ -615,6 +617,84 @@ def tbl(name: str) -> str: ), ) + # Replay-aware counter views were introduced in schema 3.0.6. + if ver >= (3, 0, 6): + # Kernel-replay dispatch counters. Dispatch 1 is replayed three times: SQ_WAVES has two + # instances in every pass, with per-pass sums 16, 18 and 20; GRBM_COUNT is collected by + # pass 0 only; SQ_INSTS_SALU by passes 1 and 2 only, with per-pass sums 5 and 6. Passes + # after the first carry replay_pass in extdata, as rocprofv3 writes them. Dispatch 2 has + # untagged SQ_WAVES rows, like a run without kernel replay. + conn.executemany( + f"INSERT INTO {tbl('rocpd_info_pmc')} " + "(id, guid, nid, pid, agent_id, target_arch, name, symbol, description, " + "component, value_type, block, expression, is_constant, is_derived, " + "spm_support, extdata) " + "VALUES (?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?,?)", + [ + ( + _PMC_GRBM_COUNT_ID, + guid, + _NID, + _PID, + _AGENT_ID, + "GPU", + "GRBM_COUNT", + "GRBM_COUNT", + "Tie High - Count Number of Clocks", + "rocm", + "ABS", + "GRBM", + "", + 0, + 0, + None, + "{}", + ), + ( + _PMC_SQ_INSTS_SALU_ID, + guid, + _NID, + _PID, + _AGENT_ID, + "GPU", + "SQ_INSTS_SALU", + "SQ_INSTS_SALU", + "Number of SALU instructions issued", + "rocm", + "ABS", + "SQ", + "", + 0, + 0, + None, + "{}", + ), + ], + ) + + def replay_pass(index): + return "{}" if index == 0 else f'{{"replay_pass":{index}}}' + + conn.executemany( + f"INSERT INTO {tbl('rocpd_pmc_event')} " + "(id, guid, event_id, pmc_id, value, extdata) VALUES (?,?,?,?,?,?)", + [ + (3, guid, _EVT_KERNEL_1, _PMC_NO_SPM_ID, 8.0, replay_pass(0)), + (4, guid, _EVT_KERNEL_1, _PMC_NO_SPM_ID, 8.0, replay_pass(0)), + (5, guid, _EVT_KERNEL_1, _PMC_GRBM_COUNT_ID, 1000.0, replay_pass(0)), + (6, guid, _EVT_KERNEL_1, _PMC_NO_SPM_ID, 9.0, replay_pass(1)), + (7, guid, _EVT_KERNEL_1, _PMC_NO_SPM_ID, 9.0, replay_pass(1)), + (8, guid, _EVT_KERNEL_1, _PMC_NO_SPM_ID, 10.0, replay_pass(2)), + (9, guid, _EVT_KERNEL_1, _PMC_NO_SPM_ID, 10.0, replay_pass(2)), + (10, guid, _EVT_KERNEL_2, _PMC_NO_SPM_ID, 4.0, "{}"), + (11, guid, _EVT_KERNEL_2, _PMC_NO_SPM_ID, 4.0, "{}"), + (12, guid, _EVT_KERNEL_1, _PMC_SQ_INSTS_SALU_ID, 2.0, replay_pass(1)), + (13, guid, _EVT_KERNEL_1, _PMC_SQ_INSTS_SALU_ID, 3.0, replay_pass(1)), + (14, guid, _EVT_KERNEL_1, _PMC_SQ_INSTS_SALU_ID, 3.0, replay_pass(2)), + (15, guid, _EVT_KERNEL_1, _PMC_SQ_INSTS_SALU_ID, 3.0, replay_pass(2)), + ], + ) + # --------------------------------------------------------------------------- # Entry point @@ -646,7 +726,31 @@ def main() -> int: schema.write_schema(conn) conn.close() + # Fail fast if the wrong schema SQL was loaded (e.g. system rocpd + stale + # ROCPD_SCHEMA_PATH). Archival schemas must not expose replay_pass; 3.0.6+ + # must. conn = sqlite3.connect(output_path) + pmc_columns = [row[1] for row in conn.execute("PRAGMA table_info(pmc_events)")] + expect_replay_pass = tuple(int(part) for part in args.version.split(".")) >= ( + 3, + 0, + 6, + ) + has_replay_pass = "replay_pass" in pmc_columns + if has_replay_pass != expect_replay_pass: + conn.close() + raise RuntimeError( + "schema {version} pmc_events columns unexpected after write_schema: " + "replay_pass present={present}, expected={expected}. columns={columns}. " + "ROCPD_SCHEMA_PATH={schema_path}".format( + version=args.version, + present=has_replay_pass, + expected=expect_replay_pass, + columns=pmc_columns, + schema_path=os.environ.get("ROCPD_SCHEMA_PATH", ""), + ) + ) + insert_minimal_data(conn, uuid, guid, args.version) conn.commit() conn.close() diff --git a/projects/rocprofiler-sdk/tests/rocpd-back-compat/validate.py b/projects/rocprofiler-sdk/tests/rocpd-back-compat/validate.py index 3b8d7a87545d..d2c5ec7f6f51 100644 --- a/projects/rocprofiler-sdk/tests/rocpd-back-compat/validate.py +++ b/projects/rocprofiler-sdk/tests/rocpd-back-compat/validate.py @@ -38,6 +38,9 @@ * CSV — graph_exec_id / graph_node_id columns absent for pre-latest schemas; out_graph_launch_trace.csv not created for pre-latest schemas. * CSV — above columns / file ARE present for the latest schema. + * CSV — counters_collection averages a counter over kernel-replay passes for the + latest schema instead of summing them, and pmc_events reports each pass. + * Perfetto — counter values for the latest schema are those pass means. * Perfetto — .pftrace exists and is non-empty for every schema. * OTF2 — "HIP Graph Launch" locations absent for pre-latest schemas; present for the latest schema (requires the ``otf2`` package). @@ -47,6 +50,7 @@ """ import csv as csv_mod +import sqlite3 import sys from pathlib import Path @@ -305,6 +309,108 @@ def test_csv_for_schema_3_0_5_changes_present(output_root, latest_schema): assert any("SE: 1" in name and "SA: 0" in name and "WGP: 3" in name for name in names) +def test_csv_counter_collection_averages_replay_passes(output_root, latest_schema): + """counters_collection must average a counter over kernel-replay passes, not sum them. + + make_db.py gives dispatch 1 SQ_WAVES in three replay passes with per-pass sums 16, 18 + and 20, GRBM_COUNT in pass 0 only (modeling later passes dedicated to another service), + and SQ_INSTS_SALU in passes 1 and 2 only with sums 5 and 6. Dispatch 2 has untagged + SQ_WAVES rows that sum to 8, modeling collection without kernel replay. + """ + counter_csv = output_root / latest_schema / "csv" / "out_counter_collection_trace.csv" + assert ( + counter_csv.exists() + ), f"out_counter_collection_trace.csv not found for schema {latest_schema}: {counter_csv}" + with open(counter_csv) as fh: + values = { + (int(row["Dispatch_Id"]), row["Counter_Name"]): float(row["Counter_Value"]) + for row in csv_mod.DictReader(fh) + } + assert values == { + (1, "SQ_WAVES"): 18.0, + (1, "GRBM_COUNT"): 1000.0, + (1, "SQ_INSTS_SALU"): 5.5, + (2, "SQ_WAVES"): 8.0, + }, f"unexpected counter_collection values for schema {latest_schema}: {values}" + + +def _open_ro_db(db_path: Path): + """Open an existing fixture DB read-only so a missing path cannot create an empty file.""" + assert db_path.is_file(), f"database not found: {db_path}" + return sqlite3.connect(f"file:{db_path.resolve().as_posix()}?mode=ro", uri=True) + + +def test_counters_collection_schema_unchanged_from_3_0_5(output_root, latest_schema): + """Replay averaging must not change the counters_collection columns.""" + + def columns(version): + db_path = output_root / version / "db" / f"schema_{version.replace('.', '_')}.db" + with _open_ro_db(db_path) as connection: + return [ + row[1] + for row in connection.execute("PRAGMA table_info(counters_collection)") + ] + + latest_columns = columns(latest_schema) + baseline_columns = columns("3.0.5") + assert latest_columns == baseline_columns, ( + f"counters_collection columns changed between 3.0.5 and {latest_schema}: " + f"3.0.5={baseline_columns} {latest_schema}={latest_columns}" + ) + + +def test_pmc_events_replay_pass_absent_before_3_0_6(output_root, old_schema): + """Archived schemas must retain their original pmc_events columns.""" + db_path = ( + output_root / old_schema / "db" / f"schema_{old_schema.replace('.', '_')}.db" + ) + with _open_ro_db(db_path) as connection: + columns = [row[1] for row in connection.execute("PRAGMA table_info(pmc_events)")] + view_sql = connection.execute( + "SELECT sql FROM sqlite_master WHERE name='pmc_events'" + ).fetchone() + assert view_sql and view_sql[0], f"pmc_events view missing for schema {old_schema}" + assert "replay_pass" not in columns, ( + f"schema {old_schema} pmc_events unexpectedly has replay_pass; " + f"columns={columns}; sql={view_sql[0]}" + ) + assert "replay_pass" not in view_sql[0], ( + f"schema {old_schema} pmc_events SQL unexpectedly mentions replay_pass: " + f"{view_sql[0]}" + ) + + +def test_pmc_events_report_replay_pass(output_root, latest_schema): + """pmc_events must report each row's kernel-replay pass, and pass 0 for untagged rows. + + SQ_WAVES of dispatch 1 is left out because the view also lists that dispatch's SPM + samples, which have no pass. + """ + db_path = ( + output_root + / latest_schema + / "db" + / f"schema_{latest_schema.replace('.', '_')}.db" + ) + with _open_ro_db(db_path) as conn: + columns = [row[1] for row in conn.execute("PRAGMA table_info(pmc_events)")] + assert ( + "replay_pass" in columns + ), f"schema {latest_schema} pmc_events missing replay_pass; columns={columns}" + rows = conn.execute( + "SELECT dispatch_id, counter_name, replay_pass, SUM(counter_value) " + "FROM pmc_events " + "WHERE counter_name IN ('GRBM_COUNT', 'SQ_INSTS_SALU') OR dispatch_id = 2 " + "GROUP BY dispatch_id, counter_name, replay_pass" + ).fetchall() + assert sorted(rows) == [ + (1, "GRBM_COUNT", 0, 1000.0), + (1, "SQ_INSTS_SALU", 1, 5.0), + (1, "SQ_INSTS_SALU", 2, 6.0), + (2, "SQ_WAVES", 0, 8.0), + ], f"unexpected pmc_events passes for schema {latest_schema}: {rows}" + + # --------------------------------------------------------------------------- # Perfetto tests — old schemas # --------------------------------------------------------------------------- @@ -329,6 +435,33 @@ def test_pftrace_completes_without_error_latest(output_root, latest_schema): assert pftrace.stat().st_size > 0, f".pftrace is empty for schema {latest_schema}" +def test_pftrace_counter_values_average_replay_passes(output_root, latest_schema): + """Perfetto counter values must be the kernel-replay pass means, not the pass sums. + + The writer annotates every kernel slice with each counter's total over all dispatches, + so only counters that make_db.py collects on a single dispatch are checked: GRBM_COUNT + (pass 0 only) and SQ_INSTS_SALU (passes 1 and 2, with sums 5 and 6). + """ + pytest.importorskip("perfetto", reason="perfetto package not installed") + from rocprofiler_sdk.pytest_utils.perfetto_reader import PerfettoReader + + pftrace = output_root / latest_schema / "pftrace" / "out_results.pftrace" + assert pftrace.exists(), f".pftrace not found for schema {latest_schema}: {pftrace}" + with PerfettoReader(str(pftrace)) as reader: + rows = reader.query_tp( + "SELECT args.key AS key, COALESCE(args.real_value, args.int_value) AS value " + "FROM slice JOIN args USING (arg_set_id) " + "WHERE args.key IN ('debug.GRBM_COUNT', 'debug.SQ_INSTS_SALU')" + ) + values = {} + for key, value in zip(rows["key"], rows["value"]): + values.setdefault(key, set()).add(float(value)) + assert values == { + "debug.GRBM_COUNT": {1000.0}, + "debug.SQ_INSTS_SALU": {5.5}, + }, f"unexpected Perfetto counter values for schema {latest_schema}: {values}" + + # --------------------------------------------------------------------------- # OTF2 tests — old schemas # --------------------------------------------------------------------------- diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/CMakeLists.txt index 6a58d8e7857f..d7a68a93e697 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/CMakeLists.txt @@ -48,8 +48,8 @@ rocprofiler_add_integration_execute_test( --pmc SQ_WAVES SQ_INSTS_VALU GRBM_GUI_ACTIVE --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_SALU --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_SMEM --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_LDS --replay-mode kernel --kernel-replay-beta-enabled - --output-format csv json -d ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay -o out -- - $ 1048576 1 + --output-format csv json rocpd -d ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay -o out + -- $ 1048576 1 DEPENDS kernel-replay TIMEOUT 120 LABELS "kernel-replay" @@ -57,14 +57,23 @@ rocprofiler_add_integration_execute_test( FIXTURES_SETUP rocprofv3-test-kernel-replay-data) # Validate replay structure: each dispatch replayed REPLAY_PASSES times, the shared sanity -# counters (SQ_WAVES, SQ_INSTS_VALU) are constant across a dispatch's passes, and the -# kernels are present. +# counters (SQ_WAVES, SQ_INSTS_VALU) are constant across a dispatch's passes, the kernels +# are present, and rocpd's counters_collection averages the shared counters over the +# passes rather than summing them. rocprofiler_add_integration_validate_test( rocprofv3-test-kernel-replay-validate BUNDLE_TESTS TEST_PATHS validate.py COPY conftest.py - ARGS --json-input ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay/out_results.json --passes - ${REPLAY_PASSES} --common-counters SQ_WAVES SQ_INSTS_VALU -x + ARGS --json-input + ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay/out_results.json + --passes + ${REPLAY_PASSES} + --common-counters + SQ_WAVES + SQ_INSTS_VALU + --rocpd-input + ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay/out_results.db + -x TIMEOUT 120 LABELS "kernel-replay" PYTHON_PATH @@ -82,8 +91,9 @@ rocprofiler_add_integration_execute_test( --pmc SQ_WAVES SQ_INSTS_VALU GRBM_GUI_ACTIVE --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_SALU --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_SMEM --pmc SQ_WAVES SQ_INSTS_VALU SQ_INSTS_LDS --replay-mode kernel --kernel-replay-beta-enabled - --output-format csv json -d ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay-interval -o - out -- $ 1048576 1 + --output-format csv json rocpd -d + ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay-interval -o out -- + $ 1048576 1 DEPENDS kernel-replay TIMEOUT 120 LABELS "kernel-replay" @@ -94,8 +104,16 @@ rocprofiler_add_integration_validate_test( rocprofv3-test-kernel-replay-interval-validate BUNDLE_TESTS TEST_PATHS validate.py COPY conftest.py - ARGS --json-input ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay-interval/out_results.json - --passes ${REPLAY_PASSES} --common-counters SQ_WAVES SQ_INSTS_VALU -x + ARGS --json-input + ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay-interval/out_results.json + --passes + ${REPLAY_PASSES} + --common-counters + SQ_WAVES + SQ_INSTS_VALU + --rocpd-input + ${CMAKE_CURRENT_BINARY_DIR}/kernel-replay-interval/out_results.db + -x TIMEOUT 120 LABELS "kernel-replay" PYTHON_PATH diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/conftest.py b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/conftest.py index b3f093b0b647..27b55f668a07 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/conftest.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/conftest.py @@ -22,6 +22,7 @@ import json import os +import sqlite3 import pytest @@ -33,6 +34,12 @@ def pytest_addoption(parser): required=True, help="rocprofv3 results JSON produced by a --replay-mode kernel --kernel-replay-beta-enabled run", ) + parser.addoption( + "--rocpd-input", + action="store", + default=None, + help="rocpd database written by the same run as --json-input", + ) parser.addoption( "--passes", action="store", @@ -74,6 +81,17 @@ def json_data(request): return json.load(f) +@pytest.fixture +def rocpd_db(request): + path = request.config.getoption("--rocpd-input") + if path is None: + pytest.skip("no --rocpd-input given") + assert os.path.isfile(path), f"missing rocpd input: {path}" + connection = sqlite3.connect(path) + yield connection + connection.close() + + @pytest.fixture def expected_passes(request): return request.config.getoption("--passes") diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/validate.py index 8b87109eff59..1dc45091f0fe 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/kernel-replay/validate.py @@ -37,9 +37,13 @@ # 5. dispatch_id identity: a replayed dispatch keeps ONE dispatch_id across its passes (only the # replay_pass index advances), and dispatch_id increments sequentially across dispatches -- one # id minted per logical dispatch, reused by every pass (never one id per pass). +# 6. With --rocpd-input, the rocpd counters_collection view reports each (dispatch, counter) as +# the mean of the passes that collected it, so the common counters are not multiplied by N, +# and every raw rocpd row carries the replay pass that pmc_events reports. # (The app verifies its own results in the generate step, guarding the restored data itself.) import collections +import json import sys # Python 3.6 compatibility: dataclasses added in 3.7 @@ -415,6 +419,77 @@ def test_common_counters_constant_across_passes(json_data, common_counter): ) +def test_rocpd_counters_collection_averages_passes(json_data, rocpd_db): + # rocpd stores every pass of a replayed dispatch under the dispatch's one event, so the + # counters_collection view must average a counter over the passes that collected it. Summing + # them reported each common counter, which every group lists, at N times its value. + per_pass = {} + for dispatch_id, entry in _records_by_dispatch(_sdk(json_data)).items(): + for batch in entry["passes"].values(): + for counter, value in batch.items(): + per_pass.setdefault((dispatch_id, counter), []).append(value) + assert any( + len(values) > 1 for values in per_pass.values() + ), "no counter was collected by more than one pass, so the run cannot show summed passes" + + actual = { + (int(dispatch_id), counter): float(value) + for dispatch_id, counter, value in rocpd_db.execute( + "SELECT dispatch_id, counter_name, value FROM counters_collection" + ) + } + assert sorted(actual) == sorted(per_pass), ( + "rocpd counters_collection and JSON disagree on the (dispatch, counter) pairs: " + f"only in rocpd {sorted(set(actual) - set(per_pass))}, " + f"only in JSON {sorted(set(per_pass) - set(actual))}" + ) + for (dispatch_id, counter), values in sorted(per_pass.items()): + mean = sum(values) / len(values) + assert actual[(dispatch_id, counter)] == pytest.approx(mean, rel=1e-9), ( + f"dispatch {dispatch_id} counter {counter}: rocpd counters_collection reports " + f"{actual[(dispatch_id, counter)]}, but its replay passes in JSON are {values} " + f"(mean {mean})" + ) + + +def test_rocpd_rows_carry_replay_pass(json_data, rocpd_db): + # The averaging above depends on the writer tagging every row: a pass after the first + # carries exactly {"replay_pass":N} in extdata and pass 0 keeps "{}". Each (dispatch, pass) + # must hold one row per instance value in that pass's JSON record, and pmc_events must + # report the same pass for each row. + want = collections.Counter() + for rec in _counter_records(_sdk(json_data)): + want[(_dispatch_id(rec), _pass_index(rec))] += len(rec.get("records", [])) + + got = collections.Counter() + for dispatch_id, extdata, rows in rocpd_db.execute( + "SELECT K.dispatch_id, E.extdata, COUNT(*) FROM rocpd_pmc_event E " + "INNER JOIN rocpd_kernel_dispatch K ON K.event_id = E.event_id AND K.guid = E.guid " + "WHERE E.sample_id IS NULL GROUP BY K.dispatch_id, E.extdata" + ): + pass_index = json.loads(extdata).get("replay_pass", 0) + tag = "{}" if pass_index == 0 else f'{{"replay_pass":{pass_index}}}' + assert ( + extdata == tag + ), f"dispatch {dispatch_id}: rocpd_pmc_event extdata {extdata!r}, expected {tag!r}" + got[(int(dispatch_id), pass_index)] += rows + assert got == want, ( + "rocpd_pmc_event rows per (dispatch, replay pass) do not match the JSON instance " + f"records: rocpd {sorted(got.items())}, JSON {sorted(want.items())}" + ) + + by_view = collections.Counter() + for dispatch_id, pass_index, rows in rocpd_db.execute( + "SELECT dispatch_id, replay_pass, COUNT(*) FROM pmc_events " + "GROUP BY dispatch_id, replay_pass" + ): + by_view[(int(dispatch_id), int(pass_index))] += rows + assert by_view == want, ( + "pmc_events rows per (dispatch, replay_pass) do not match the JSON instance records: " + f"rocpd {sorted(by_view.items())}, JSON {sorted(want.items())}" + ) + + def test_each_pass_collects_distinct_batch(json_data, expected_passes, common_counters): # Metric 3: each pass collects exactly one unique (non-common) counter, and those per-pass # counters differ — proving real multi-pass batch rotation rather than one pass carrying