diff --git a/projects/rocprofiler-sdk/CHANGELOG.md b/projects/rocprofiler-sdk/CHANGELOG.md index 799086221d95..e8c5905f24d9 100644 --- a/projects/rocprofiler-sdk/CHANGELOG.md +++ b/projects/rocprofiler-sdk/CHANGELOG.md @@ -10,6 +10,10 @@ Full documentation for ROCprofiler-SDK is available at [rocm.docs.amd.com/projec ### Resolved issues +**rocpd:** + + - `counters_collection` no longer sums a counter that kernel replay collects once per `--pmc` group when every pass reports the same value. Three readings of 512 report 512. Readings that differ, such as hardware instances of one counter, are still summed. The individual readings remain in `rocpd_pmc_event`. + ### Known issues ### Removed diff --git a/projects/rocprofiler-sdk/source/docs/_toc.yml.in b/projects/rocprofiler-sdk/source/docs/_toc.yml.in index 4767a7c6c69b..4eec10a526d2 100644 --- a/projects/rocprofiler-sdk/source/docs/_toc.yml.in +++ b/projects/rocprofiler-sdk/source/docs/_toc.yml.in @@ -126,6 +126,8 @@ subtrees: title: Memory snapshot and restore - file: conceptual/kernel_replay/kernel_replay_performance title: Performance assessment + - file: conceptual/kernel_replay/kernel_replay_rocpd_schema + title: rocpd schema - file: conceptual/kernel_replay/kernel_replay_testing title: Test coverage - file: conceptual/kernel_replay/kernel_replay_downstream_tools diff --git a/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/index.md b/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/index.md index 80d392e73e0f..a0dd99e79077 100644 --- a/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/index.md +++ b/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/index.md @@ -73,6 +73,7 @@ same domain for timing, PC sampling, or thread trace. per-dispatch snapshot cost model, the range the regression tests actually cover, the specific performance problems visible in the implementation, and how the cost behaves as device capacity outpaces host-link bandwidth across accelerator generations. +- **[rocpd schema](kernel_replay_rocpd_schema.md)** — `counters_collection` reports a repeated counter once when every replay pass reads the same value. Readings that differ are still summed. The individual readings remain in `rocpd_pmc_event`. - **[Test coverage](kernel_replay_testing.md)** — what replay is tested for at each level, which checks need a GPU and which deliberately do not, the known gaps, and where a new test belongs. - **[Downstream tools](kernel_replay_downstream_tools.md)** — what it would take for diff --git a/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/kernel_replay_rocpd_schema.md b/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/kernel_replay_rocpd_schema.md new file mode 100644 index 000000000000..b77b5efeab76 --- /dev/null +++ b/projects/rocprofiler-sdk/source/docs/conceptual/kernel_replay/kernel_replay_rocpd_schema.md @@ -0,0 +1,26 @@ +# Kernel replay in the rocpd schema + +Kernel replay executes one dispatch once per counter group. Every pass reports the same +`dispatch_id`. rocpd stores every pass's counter readings as `rocpd_pmc_event` rows on that +dispatch's event. `counters_collection` used to `SUM` those rows, so a sanity counter repeated in +three groups was reported three times too high. + +```bash +rocprofv3 --pmc SQ_WAVES,SQ_INSTS_VALU,GRBM_COUNT \ + --pmc SQ_WAVES,SQ_INSTS_VALU,GRBM_GUI_ACTIVE \ + --pmc SQ_WAVES,SQ_INSTS_VALU,SQ_INSTS_SALU \ + --replay-mode kernel --kernel-replay-beta-enabled -- ./app +``` + +On a kernel that launches 512 waves, each pass reads `SQ_WAVES` as 512. The view used to report +1536. Identical readings now report that value once, so this reports 512. Readings that are not +identical, such as hardware instances of one counter, are still summed. A counter collected once +is unchanged. A run that never replayed is unchanged when each counter name is collected once. + +The three readings stay in `rocpd_pmc_event`. A query against that table still lists 512, 512, and +512. For a deterministic counter those readings are the same number, so the view reports +that number instead of adding the copies. + +`counters_collection` stays one row per dispatch per counter. That is the file most people read. +Keeping a row per pass, with its own timestamps, would be a larger schema change, and it is not +what this view is for. One default cannot serve both. 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..590b62541cea 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 @@ -784,7 +784,13 @@ SELECT PMC_I.expression, PMC_I.value_type, PMC_I.id AS counter_id, - SUM(PMC_E.value) AS value, + -- Several rows can share a counter name. Hardware instances of one collection differ and + -- still add up. Kernel replay writes the same reading again for each pass, and those + -- copies are one measurement, so identical readings report that value once. + CASE + WHEN MIN(PMC_E.value) = MAX(PMC_E.value) THEN MIN(PMC_E.value) + ELSE SUM(PMC_E.value) + END AS value, K.start, K.end, PMC_I.is_constant, diff --git a/projects/rocprofiler-sdk/tests/CMakeLists.txt b/projects/rocprofiler-sdk/tests/CMakeLists.txt index d8d4ac813828..30e690efeb03 100644 --- a/projects/rocprofiler-sdk/tests/CMakeLists.txt +++ b/projects/rocprofiler-sdk/tests/CMakeLists.txt @@ -117,6 +117,7 @@ add_subdirectory(environment) add_subdirectory(rocpd) add_subdirectory(rocpd-api) add_subdirectory(rocpd-back-compat) +add_subdirectory(rocpd-kernel-replay-schema) # rocattach validation tests add_subdirectory(rocattach) diff --git a/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/CMakeLists.txt b/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/CMakeLists.txt new file mode 100644 index 000000000000..32309aee5db7 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/CMakeLists.txt @@ -0,0 +1,34 @@ +# +# rocpd kernel replay schema tests: drive the shipped SQL against SQLite, no GPU +# required +# +cmake_minimum_required(VERSION 3.21.0 FATAL_ERROR) + +project( + rocprofiler-sdk-tests-rocpd-kernel-replay-schema + LANGUAGES NONE + VERSION 0.0.0) + +find_package(Python3 REQUIRED COMPONENTS Interpreter) + +set(rocpd-kernel-replay-schema-path "${CMAKE_BINARY_DIR}/share/rocprofiler-sdk-rocpd") + +# The tests build their databases from the schema SQL directly, so there is nothing to +# profile and no application to run. The execute step only exists to provide the fixture +# that rocprofiler_add_integration_validate_test requires. +rocprofiler_add_integration_execute_test( + rocpd-kernel-replay-schema-test + COMMAND ${Python3_EXECUTABLE} -c "pass" + DEPENDS rocprofiler-sdk::rocprofv3 + TIMEOUT 10 + LABELS "integration-tests;rocpd;kernel-replay" + FIXTURES_SETUP rocpd-kernel-replay-schema) + +rocprofiler_add_integration_validate_test( + rocpd-kernel-replay-schema-test + TEST_PATHS validate.py + CONFIG pytest.ini + TIMEOUT 60 + LABELS "integration-tests;rocpd;kernel-replay" + ENVIRONMENT "ROCPD_SCHEMA_PATH=${rocpd-kernel-replay-schema-path}" + FIXTURES_REQUIRED rocpd-kernel-replay-schema) diff --git a/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/pytest.ini b/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/pytest.ini new file mode 100644 index 000000000000..8bf72b0989bc --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/pytest.ini @@ -0,0 +1,5 @@ + +[pytest] +addopts = --durations=20 -rA -s +testpaths = validate.py +pythonpath = @ROCPROFILER_SDK_TESTS_BINARY_DIR@/pytest-packages diff --git a/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/validate.py b/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/validate.py new file mode 100644 index 000000000000..e555e490e7f8 --- /dev/null +++ b/projects/rocprofiler-sdk/tests/rocpd-kernel-replay-schema/validate.py @@ -0,0 +1,306 @@ +#!/usr/bin/env python3 + +# MIT License +# +# Copyright (c) 2024-2026 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. + +"""Schema-level checks for kernel replay counter values in rocpd. + +Kernel replay executes one dispatch once per counter group and every pass reports the same +dispatch_id. The counter readings are stored, and counters_collection reports their average so a +sanity counter repeated in every group is not multiplied by the pass count. These tests drive the +shipped SQL against SQLite, with no GPU and no profiling run. +""" + +import os +import sqlite3 +from pathlib import Path + +import pytest + +# SQ_WAVES is repeated in every group. Each pass reports 512 waves. +REPLAY_PASSES = 3 +SQ_WAVES_PER_PASS = 512.0 + +# One counter unique to each group. Those were never multiplied. +GROUP_UNIQUE_COUNTERS = ["GRBM_COUNT", "GRBM_GUI_ACTIVE", "SQ_INSTS_SALU"] + + +def schema_dir() -> Path: + """Directory holding the current schema SQL, from the build tree or the source tree.""" + + env_path = os.environ.get("ROCPD_SCHEMA_PATH") + if env_path and (Path(env_path) / "rocpd_tables.sql").exists(): + return Path(env_path) + + source_dir = ( + Path(__file__).resolve().parents[2] / "source" / "share" / "rocprofiler-sdk-rocpd" + ) + if (source_dir / "rocpd_tables.sql").exists(): + return source_dir + + pytest.skip("could not locate the rocpd schema SQL") + + +def build_database(version_dir: Path) -> sqlite3.Connection: + """Create an in-memory database from the schema SQL in the given directory.""" + + conn = sqlite3.connect(":memory:") + for filename in ("rocpd_tables.sql", "rocpd_views.sql", "data_views.sql"): + sql = (version_dir / filename).read_text() + # The shipped SQL is a template: {{uuid}} suffixes the physical table names for merged + # databases and {{guid}} tags the rows with their source. A single unmerged database uses + # an empty uuid, which is what rocprofv3 itself writes. + sql = sql.replace("{{uuid}}", "").replace("{{guid}}", "test-guid") + conn.executescript(sql) + return conn + + +def insert_common_rows(conn: sqlite3.Connection, counter_names) -> None: + """Insert the node/process/agent/kernel rows every counter row has to join against.""" + + cur = conn.cursor() + cur.execute( + "INSERT INTO rocpd_info_node (id, guid, hash, machine_id, hostname)" + " VALUES (1, 'test-guid', 1, 'machine', 'host')" + ) + cur.execute( + "INSERT INTO rocpd_info_process (id, guid, nid, pid) VALUES (1, 'test-guid', 1, 1000)" + ) + cur.execute( + "INSERT INTO rocpd_info_thread (id, guid, nid, pid, tid)" + " VALUES (1, 'test-guid', 1, 1, 2000)" + ) + cur.execute( + "INSERT INTO rocpd_info_agent" + " (id, guid, nid, pid, type, absolute_index, logical_index, type_index, name)" + " VALUES (1, 'test-guid', 1, 1, 'GPU', 0, 0, 0, 'agent-0')" + ) + cur.execute( + "INSERT INTO rocpd_info_queue (id, guid, nid, pid, name)" + " VALUES (1, 'test-guid', 1, 1, 'queue-0')" + ) + cur.execute( + "INSERT INTO rocpd_info_stream (id, guid, nid, pid, name)" + " VALUES (1, 'test-guid', 1, 1, 'stream-0')" + ) + cur.execute( + "INSERT INTO rocpd_info_code_object (id, guid, nid, pid, agent_id, storage_type)" + " VALUES (1, 'test-guid', 1, 1, 1, 'FILE')" + ) + cur.execute( + "INSERT INTO rocpd_info_kernel_symbol" + " (id, guid, nid, pid, code_object_id, kernel_name, display_name," + " sgpr_count, arch_vgpr_count, accum_vgpr_count)" + " VALUES (1, 'test-guid', 1, 1, 1, 'saxpy', 'saxpy', 16, 8, 4)" + ) + cur.execute( + "INSERT INTO rocpd_string (id, guid, string) VALUES (1, 'test-guid', 'KERNEL')" + ) + + for pmc_id, name in enumerate(counter_names, start=1): + cur.execute( + "INSERT INTO rocpd_info_pmc" + " (id, guid, nid, pid, agent_id, target_arch, name, symbol, value_type," + " block, component, is_constant, is_derived)" + " VALUES (?, 'test-guid', 1, 1, 1, 'GPU', ?, ?, 'ABS', 'SQ', 'gfx', 0, 0)", + (pmc_id, name, name), + ) + conn.commit() + + +def insert_dispatch( + conn: sqlite3.Connection, *, row_id: int, event_id: int, dispatch_id: int, start: int +) -> None: + """Insert one dispatch and the event row the counter rows attach to.""" + + cur = conn.cursor() + cur.execute( + "INSERT INTO rocpd_event (id, guid, category_id, stack_id, parent_stack_id, correlation_id)" + " VALUES (?, 'test-guid', 1, ?, ?, ?)", + (event_id, dispatch_id, dispatch_id, dispatch_id), + ) + cur.execute( + "INSERT INTO rocpd_kernel_dispatch (" + " id, guid, nid, pid, tid, agent_id, kernel_id, dispatch_id, queue_id, stream_id," + " start, end, workgroup_size_x, workgroup_size_y, workgroup_size_z," + " grid_size_x, grid_size_y, grid_size_z, event_id" + ") VALUES (?, 'test-guid', 1, 1, 1, 1, 1, ?, 1, 1, ?, ?, 64, 1, 1, 512, 1, 1, ?)", + (row_id, dispatch_id, start, start + 100, event_id), + ) + conn.commit() + + +def insert_counter( + conn: sqlite3.Connection, *, row_id: int, event_id: int, pmc_id: int, value: float +) -> None: + conn.execute( + "INSERT INTO rocpd_pmc_event (id, guid, event_id, pmc_id, value)" + " VALUES (?, 'test-guid', ?, ?, ?)", + (row_id, event_id, pmc_id, value), + ) + conn.commit() + + +def populate_replay_run(conn: sqlite3.Connection) -> None: + """One dispatch, three passes, all counter rows on the first event. + + That is the shape the export writes: later passes do not get their own dispatch row, and + every reading is attached to the event recorded for that dispatch_id. + """ + + counter_names = ["SQ_WAVES"] + GROUP_UNIQUE_COUNTERS + insert_common_rows(conn, counter_names) + insert_dispatch(conn, row_id=1, event_id=1, dispatch_id=1, start=1000) + + pmc_ids = {name: idx for idx, name in enumerate(counter_names, start=1)} + counter_row_id = 1 + for pass_index in range(REPLAY_PASSES): + insert_counter( + conn, + row_id=counter_row_id, + event_id=1, + pmc_id=pmc_ids["SQ_WAVES"], + value=SQ_WAVES_PER_PASS, + ) + counter_row_id += 1 + insert_counter( + conn, + row_id=counter_row_id, + event_id=1, + pmc_id=pmc_ids[GROUP_UNIQUE_COUNTERS[pass_index]], + value=100.0 * (pass_index + 1), + ) + counter_row_id += 1 + + +def populate_plain_run(conn: sqlite3.Connection) -> None: + """Three ordinary dispatches, one reading of each counter. AVG must equal that reading.""" + + counter_names = ["SQ_WAVES", "SQ_INSTS_VALU"] + insert_common_rows(conn, counter_names) + counter_row_id = 1 + for dispatch_id in range(1, 4): + insert_dispatch( + conn, + row_id=dispatch_id, + event_id=dispatch_id, + dispatch_id=dispatch_id, + start=1000 * dispatch_id, + ) + for pmc_id, value in ((1, 512.0 * dispatch_id), (2, 64.0 * dispatch_id)): + insert_counter( + conn, + row_id=counter_row_id, + event_id=dispatch_id, + pmc_id=pmc_id, + value=value, + ) + counter_row_id += 1 + + +@pytest.fixture +def current_schema(): + return schema_dir() + + +def test_repeated_counter_is_averaged(current_schema): + """SQ_WAVES in all three groups reads 512 once, not 1536.""" + + conn = build_database(current_schema) + populate_replay_run(conn) + rows = conn.execute( + "SELECT value FROM counters_collection WHERE counter_name = 'SQ_WAVES'" + ).fetchall() + assert len(rows) == 1 + assert rows[0][0] == SQ_WAVES_PER_PASS + + +def test_raw_readings_remain(current_schema): + """The view averages. It does not delete the per-pass readings.""" + + conn = build_database(current_schema) + populate_replay_run(conn) + rows = conn.execute( + "SELECT PMC_E.value FROM rocpd_pmc_event PMC_E" + " JOIN rocpd_info_pmc PMC_I ON PMC_I.id = PMC_E.pmc_id" + " WHERE PMC_I.name = 'SQ_WAVES' ORDER BY PMC_E.id" + ).fetchall() + assert [row[0] for row in rows] == [SQ_WAVES_PER_PASS] * REPLAY_PASSES + + +def test_sum_of_the_same_rows_is_the_old_inflated_value(current_schema): + """Documents the number SUM used to report for this shape.""" + + conn = build_database(current_schema) + populate_replay_run(conn) + total = conn.execute( + "SELECT SUM(PMC_E.value) FROM rocpd_pmc_event PMC_E" + " JOIN rocpd_info_pmc PMC_I ON PMC_I.id = PMC_E.pmc_id" + " WHERE PMC_I.name = 'SQ_WAVES'" + ).fetchone()[0] + assert total == SQ_WAVES_PER_PASS * REPLAY_PASSES + + +def test_group_unique_counters_are_unaffected(current_schema): + """A counter in only one group is collected once, so AVG does not change it.""" + + conn = build_database(current_schema) + populate_replay_run(conn) + for pass_index, name in enumerate(GROUP_UNIQUE_COUNTERS): + rows = conn.execute( + "SELECT value FROM counters_collection WHERE counter_name = ?", + (name,), + ).fetchall() + assert rows == [(100.0 * (pass_index + 1),)] + + +def test_different_readings_of_one_counter_still_sum(current_schema): + """Hardware instances of one counter differ, so the view still adds them.""" + + conn = build_database(current_schema) + insert_common_rows(conn, ["SQ_WAVES"]) + insert_dispatch(conn, row_id=1, event_id=1, dispatch_id=1, start=1000) + insert_counter(conn, row_id=1, event_id=1, pmc_id=1, value=100.0) + insert_counter(conn, row_id=2, event_id=1, pmc_id=1, value=300.0) + rows = conn.execute( + "SELECT value FROM counters_collection WHERE counter_name = 'SQ_WAVES'" + ).fetchall() + assert rows == [(400.0,)] + + +def test_single_reading_is_unchanged(current_schema): + """A dispatch with one reading of a counter reports that reading.""" + + conn = build_database(current_schema) + populate_plain_run(conn) + rows = conn.execute( + "SELECT dispatch_id, counter_name, value FROM counters_collection" + " ORDER BY dispatch_id, counter_name" + ).fetchall() + assert rows == [ + (1, "SQ_INSTS_VALU", 64.0), + (1, "SQ_WAVES", 512.0), + (2, "SQ_INSTS_VALU", 128.0), + (2, "SQ_WAVES", 1024.0), + (3, "SQ_INSTS_VALU", 192.0), + (3, "SQ_WAVES", 1536.0), + ] diff --git a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py index c76f14914bd2..4c6099a80a07 100644 --- a/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py +++ b/projects/rocprofiler-sdk/tests/rocprofv3/rocpd/validate.py @@ -77,11 +77,19 @@ def test_csv_data(csv_data, json_data): def test_counter_collection_fields_match_across_rocpd_csv_json( csv_data, json_data, db_input ): + def combine(values): + # Same rule as counters_collection: identical readings are one measurement + # (a counter repeated across kernel replay passes). Readings that differ + # are hardware instances and still sum. + if min(values) == max(values): + return values[0] + return sum(values) + def aggregate(records): - result = defaultdict(float) + grouped = defaultdict(list) for dispatch_id, counter_name, value in records: - result[(int(dispatch_id), counter_name)] += float(value) - return dict(result) + grouped[(int(dispatch_id), counter_name)].append(float(value)) + return {key: combine(values) for key, values in grouped.items()} data = json_data["rocprofiler-sdk-tool"] counter_names = {