From 33e78837eb928296df4c057e5112758c3fadebe6 Mon Sep 17 00:00:00 2001 From: Vignesh Edithal Date: Wed, 16 Sep 2026 13:27:38 +0000 Subject: [PATCH 01/45] [rocprofiler-compute] Open Q3 Sprint 4 (Sep 14 - Sep 28) Empty commit to open the sprint integration PR against develop. Co-Authored-By: Claude Opus 5 (1M context) From 6230b451f868d38c4bbee32ecfbdd7b0198bcdb2 Mon Sep 17 00:00:00 2001 From: Vignesh Edithal Date: Sun, 16 Aug 2026 07:17:11 +0000 Subject: [PATCH 02/45] [rocprofiler-compute] Open Q3 Sprint 3 (Aug 17 - Sep 11) Empty commit to open the sprint integration PR against develop. Co-Authored-By: Claude Opus 5 (1M context) From 2e57a89783e20da8fe4fe57c08fb0beb1b5bb024 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Mon, 21 Sep 2026 13:34:13 -0500 Subject: [PATCH 03/45] script for analysis run --- .../run_rocprof_analysis.sh | 30 +++++++++++++++++++ 1 file changed, 30 insertions(+) create mode 100755 projects/rocprofiler-compute/run_rocprof_analysis.sh diff --git a/projects/rocprofiler-compute/run_rocprof_analysis.sh b/projects/rocprofiler-compute/run_rocprof_analysis.sh new file mode 100755 index 000000000000..b87e9bb3fb68 --- /dev/null +++ b/projects/rocprofiler-compute/run_rocprof_analysis.sh @@ -0,0 +1,30 @@ +#!/bin/bash + +set -euo pipefail + +if [ $# -ne 1 ]; then + echo "Usage: $0 " + exit 1 +fi + +FOLDER="$1" + +if [ ! -d "$FOLDER" ]; then + echo "Error: '$FOLDER' is not a directory" + exit 1 +fi + +for full_path in "$FOLDER"/*; do + [ -f "$full_path" ] || continue + + file_name="$(basename "$full_path")" + + echo "=== Profiling: $file_name ===" + ./src/rocprof-compute profile --roof-only \ + --output-directory "/work-space/output/bmm/$file_name" \ + --overwrite \ + -- python "$full_path" + + echo "=== Analyzing: $file_name ===" + ./src/rocprof-compute analyze -p "/work-space/output/bmm/$file_name" +done From 0d8390cd5d6cd2e20690338a90b24e4e9a787119 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Mon, 21 Sep 2026 13:34:24 -0500 Subject: [PATCH 04/45] benchmark fixes --- .../src/roofline/benchmark/benchmark_base.py | 9 +++++++-- .../src/roofline/benchmark/gfx9/benchmark_gfx950.py | 10 +++++----- 2 files changed, 12 insertions(+), 7 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py index 302bd9009b59..08bc3b3d2743 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py @@ -43,7 +43,9 @@ PerfMetrics = namedtuple("PerfMetrics", ["mean", "low", "high"]) DEFAULT_WORKGROUP_SIZE = 256 -DEFAULT_NUM_EXPERIMENTS = 100 +DEFAULT_WORKGROUPS = 8192 +DEFAULT_THREADS = DEFAULT_WORKGROUP_SIZE * DEFAULT_WORKGROUPS +DEFAULT_NUM_EXPERIMENTS = 10 DEFAULT_NUM_ITERS = 10 # Number of FMA operations per thread iteration in VALU benchmark. @@ -69,6 +71,8 @@ def __init__(self, device_id: int, cache_sizes: dict) -> None: self.csv_cols_map: dict[str, str] self.WAVEFRONT_SIZE: int self.MATRIX_OPS_TYPE: str + self.event_start = hip.hipEventCreate() + self.event_stop = hip.hipEventCreate() self.cache_sizes = cache_sizes @@ -168,6 +172,7 @@ def show_progress(self, pct: float) -> None: def calc_stats(self, samples: list) -> Stats: """Returns a named tuple with the mean, std deviation and confidence.""" mean = sum(samples) / len(samples) + m = max(samples) stdev = 0.0 @@ -176,7 +181,7 @@ def calc_stats(self, samples: list) -> Stats: stdev = math.sqrt(stdev / len(samples)) - return Stats(mean, stdev, 1.96 * stdev / math.sqrt(len(samples))) + return Stats(m, stdev, 1.96 * stdev / math.sqrt(len(samples))) class Program: """Helper class for loading and compiling kernels.""" diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py b/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py index a7a9351a5853..49337c2d4cbb 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py @@ -27,7 +27,7 @@ def __init__(self, device_id: int, cache_sizes: dict) -> None: "F6F4": 131072, # Mixed precision F6 x F4 "MXF8": 131072, # microscaling FP8 format "F8": 32768, - "F16": 32768, + "F16": 16384, "F32": 4096, "BF16": 32768, "I8": 65536, @@ -77,13 +77,13 @@ def set_kernel_source(self) -> None: + """ extern "C" __global__ void mfma_f16(int iter, float *dummy) { - vec16 result = {0}; + vec4 result = {0}; vec8<__fp16> a; a[7] = a[6] = a[5] = a[4] = a[3] = a[2] = a[1] = a[0] = threadIdx.x; for(int i = 0; i < iter; ++i) { - result = __builtin_amdgcn_mfma_f32_32x32x16_f16(a, a, result, 0, 0, 0); + result = __builtin_amdgcn_mfma_f32_16x16x32_f16(a, a, result, 0, 0, 0); } if (result[0] != 2*result[0]) @@ -100,12 +100,12 @@ def set_kernel_source(self) -> None: extern "C" __global__ void mfma_bf16(int iter, float *dummy) { vec16 result = {0}; - vec8 a; + vec8<__fp16> a; a[7] = a[6] = a[5] = a[4] = a[3] = a[2] = a[1] = a[0] = threadIdx.x; for(int i = 0; i < iter; ++i) { - result = __builtin_amdgcn_mfma_f32_32x32x16_bf16(a, a, result, 0, 0, 0); + result = __builtin_amdgcn_mfma_f32_32x32x16_f16(a, a, result, 0, 0, 0); } if (result[0] != 2*result[0]) From 596fe87bbe6edf908aeb31640c54f2d0200cedd2 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Tue, 22 Sep 2026 07:52:31 -0500 Subject: [PATCH 05/45] Adding bandwidth to roofline hover --- projects/rocprofiler-compute/src/roofline/roofline_hover.py | 5 +++-- projects/rocprofiler-compute/src/roofline/roofline_main.py | 6 +++++- 2 files changed, 8 insertions(+), 3 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index e21edf49e671..53805ac30b47 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -49,6 +49,7 @@ def build_kernel_hover_template( f"Achieved throughput: %{{y:,.3f}} {unit}", f"Peak throughput: %{{customdata[0]}} {unit}", "Percent of roofline achieved: %{customdata[1]} %", + "Cache level bandwidth: %{customdata[2]}", f"Performance limiter: {limiter}", f"Total dispatches: {_format_integer(count)}", f"Aggregate time in kernel: {time_txt}", @@ -67,7 +68,7 @@ def build_roof_hover( roof this slope caps against, each labeled with its datatype.""" rows = [ "Model: throughput = min(bandwidth \u00d7 AI, compute peak).", - f"Bandwidth (slope): {_format_bandwidth(bandwidth)}", + f"Bandwidth (slope): {format_bandwidth(bandwidth)}", ] if compute_peaks: rows.append("Compute peaks (flat roofs):") @@ -116,7 +117,7 @@ def _hover(header: str, rows: list[str]) -> str: return "
".join([header, "", *rows]) + "" -def _format_bandwidth(gb_per_s: float) -> str: +def format_bandwidth(gb_per_s: float) -> str: """Bandwidth as GB/s, switching to TB/s at >= 1000 GB/s so the roof hover stays readable.""" try: diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 3070a1f573bd..e66d7c6a538f 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -17,6 +17,7 @@ build_compute_peak_hover, build_kernel_hover_template, build_roof_hover, + format_bandwidth, format_hover_number, wrap_hover_name, ) @@ -489,9 +490,11 @@ def _build_kernel_points( if not (ai_value > 0 and performance > 0): continue + cache_key = cache_level.removeprefix("ai_") + bandwidth = self._peak_value(ceiling_data, cache_key) roof_perf = self._roof_value_at( ai_value=ai_value, - cache_key=cache_level.removeprefix("ai_"), + cache_key=cache_key, ceiling_data=ceiling_data, cap=compute_cap, ) @@ -503,6 +506,7 @@ def _build_kernel_points( "hoverCells": [ format_hover_number(roof_perf, ",.3f"), format_hover_number(pct_roof, ".4f"), + format_bandwidth(bandwidth) if bandwidth is not None else "N/A", ], }) level_ai[level_name] = ai_value From 8ed7550523bcca0960eb7f03b3081edf7d0c39fd Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Tue, 22 Sep 2026 19:59:31 +0100 Subject: [PATCH 06/45] roofline exported as csv --- .../rocprof_compute_analyze/analysis_cli.py | 1 + .../src/roofline/roofline_csv.py | 69 +++++++++++ .../src/roofline/roofline_main.py | 30 +++++ .../integration/test_roofline_analyze.py | 33 ++++++ .../tests/unit/roofline/test_roofline_csv.py | 111 ++++++++++++++++++ .../tests/unit/roofline/test_roofline_main.py | 38 ++++++ 6 files changed, 282 insertions(+) create mode 100644 projects/rocprofiler-compute/src/roofline/roofline_csv.py create mode 100644 projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py diff --git a/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py b/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py index 4d882c0ecb16..c657c61552d2 100644 --- a/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py +++ b/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py @@ -244,6 +244,7 @@ def run_analysis(self) -> None: flops_dt, ) = roof_obj.construct_plotly_figures(ai_data=ai_data) roof_obj.save_html_files(ops_fig, flops_fig, ops_dt, flops_dt) + roof_obj.save_kernel_csv() else: console_warning( "roofline", diff --git a/projects/rocprofiler-compute/src/roofline/roofline_csv.py b/projects/rocprofiler-compute/src/roofline/roofline_csv.py new file mode 100644 index 000000000000..3e6c52136b54 --- /dev/null +++ b/projects/rocprofiler-compute/src/roofline/roofline_csv.py @@ -0,0 +1,69 @@ +# Copyright (c) Advanced Micro Devices, Inc. +# SPDX-License-Identifier: MIT + +"""Per-kernel roofline data export, alongside the interactive roofline HTML.""" + +from pathlib import Path +from typing import Any, Union + +import pandas as pd + +from utils.logger import console_log + +KERNEL_ROOFLINE_CSV_FILENAME = "roofline_kernels.csv" + +KERNEL_ROOFLINE_CSV_COLUMNS = [ + "Kernel_Name", + "Cache_Level", + "Arithmetic_Intensity", + "Performance_GFLOPs", + "Roofline_Peak_GFLOPs", + "Pct_of_Roofline", + "Bandwidth_GBps", + "Count", + "Total_Time", + "Time_Unit", + "Pct_Runtime", +] + + +def collect_kernel_roofline_rows( + kernels: list[dict[str, Any]], time_unit: str +) -> list[dict[str, Any]]: + """Flatten the FLOP figure's per-kernel view-model data into tidy rows. + + One row per kernel per memory level actually plotted, read straight from + RooflineViewModel.kernels (built by Roofline._build_kernel_traces), so + these numbers can never diverge from what the interactive HTML shows. + """ + rows: list[dict[str, Any]] = [] + for kernel in kernels: + for point in kernel["points"]: + rows.append({ + "Kernel_Name": kernel["name"], + "Cache_Level": point["peak"], + "Arithmetic_Intensity": point["ai"], + "Performance_GFLOPs": point["perf"], + "Roofline_Peak_GFLOPs": point["roofPerf"], + "Pct_of_Roofline": point["pctRoof"], + "Bandwidth_GBps": point["bandwidth"], + "Count": kernel.get("count"), + "Total_Time": kernel.get("totalTime"), + "Time_Unit": time_unit, + "Pct_Runtime": kernel.get("pctRuntime"), + }) + return rows + + +def build_kernel_roofline_dataframe(rows: list[dict[str, Any]]) -> pd.DataFrame: + """Assemble tidy per-kernel roofline rows into a DataFrame with a fixed + column order.""" + return pd.DataFrame(rows, columns=KERNEL_ROOFLINE_CSV_COLUMNS) + + +def write_kernel_roofline_csv(df: pd.DataFrame, workload_dir: Union[str, Path]) -> Path: + """Write the per-kernel roofline DataFrame to roofline_kernels.csv.""" + output_file = Path(workload_dir) / KERNEL_ROOFLINE_CSV_FILENAME + df.to_csv(output_file, index=False) + console_log("roofline", f"Saved per-kernel roofline data to {output_file}") + return output_file diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index e66d7c6a538f..793befd94d1d 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -12,6 +12,11 @@ import plotly.colors as pcolors import plotly.graph_objects as go +from roofline.roofline_csv import ( + build_kernel_roofline_dataframe, + collect_kernel_roofline_rows, + write_kernel_roofline_csv, +) from roofline.roofline_frame import canonical_frame from roofline.roofline_hover import ( build_compute_peak_hover, @@ -458,6 +463,8 @@ def _build_kernel_traces( "color": color, "points": points, "pctRuntime": pct_val, + "count": count_val, + "totalTime": time_val, }) return traces, kernels_model @@ -503,6 +510,9 @@ def _build_kernel_points( "peak": level_name, "ai": ai_value, "perf": performance, + "roofPerf": roof_perf, + "pctRoof": pct_roof, + "bandwidth": bandwidth, "hoverCells": [ format_hover_number(roof_perf, ",.3f"), format_hover_number(pct_roof, ".4f"), @@ -613,6 +623,26 @@ def save_html_files( Path(path).write_text(document, encoding="utf-8") console_log("roofline", "Roofline HTML files saved.") + def save_kernel_csv(self) -> Optional[Path]: + """Write the per-kernel roofline data next to the HTML, as a tidy CSV. + + Reuses the FLOP figure's kernel view-model, so the CSV and the + interactive HTML always agree. Call after construct_plotly_figures, + which is what populates that view-model. + """ + view_model = self.__view_models.get("FLOP") + if view_model is None or not view_model.kernels: + return None + + time_unit = (self.__ai_data or {}).get("timeUnit", "") + rows = collect_kernel_roofline_rows(view_model.kernels, time_unit) + if not rows: + return None + + df = build_kernel_roofline_dataframe(rows) + workload_dir = self.__run_parameters["workload_dir"] + return write_kernel_roofline_csv(df, workload_dir) + def _combined_html_figure( self, ops_figure: Optional[go.Figure], diff --git a/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py b/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py index 137624654bbd..3c19451a8bc0 100644 --- a/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py +++ b/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py @@ -17,6 +17,7 @@ import common import pytest +from roofline.roofline_csv import KERNEL_ROOFLINE_CSV_FILENAME from tests.integration import common as integration_common config = {} @@ -74,6 +75,38 @@ def test_analyze_generates_roofline_html( common.clean_output_dir(config["cleanup"], workload_dir) +def test_analyze_generates_kernel_roofline_csv( + binary_handler_analyze_rocprof_compute: Callable[[list[str]], int], +) -> None: + """ + Analyze writes roofline_kernels.csv alongside the roofline HTML, with the + same per-kernel data plotted in the interactive figure. + + The MI200 fixture's kernel is a pure memory-bandwidth microbenchmark with + zero measured FLOPs, so it draws no dot on the FLOP figure and no CSV row + either -- matching the empty-HTML-plot case exactly, and confirming the + export doesn't write a spurious empty file. Coverage for the populated + (non-empty rows, expected columns) case lives in the unit tests, where + kernel performance data can be constructed directly. + """ + workload_dir = integration_common.setup_workload_dir(roofline_dir) + + assert (Path(workload_dir) / "roofline.csv").exists() + + code = binary_handler_analyze_rocprof_compute([ + "analyze", + "--path", + workload_dir, + "--roofline-data-type", + "FP32", + ]) + assert code == 0 + + assert not (Path(workload_dir) / KERNEL_ROOFLINE_CSV_FILENAME).exists() + + common.clean_output_dir(config["cleanup"], workload_dir) + + def test_analyze_roofline_datatype_independently( binary_handler_analyze_rocprof_compute: Callable[[list[str]], int], ) -> None: diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py new file mode 100644 index 000000000000..b60099abaf01 --- /dev/null +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py @@ -0,0 +1,111 @@ +# Copyright (c) Advanced Micro Devices, Inc. +# SPDX-License-Identifier: MIT + +"""Unit coverage for the per-kernel roofline CSV export: flattening the FLOP +figure's kernel view-model into tidy rows and writing them to disk. +""" + +from pathlib import Path + +import pandas as pd + +from roofline.roofline_csv import ( + KERNEL_ROOFLINE_CSV_COLUMNS, + KERNEL_ROOFLINE_CSV_FILENAME, + build_kernel_roofline_dataframe, + collect_kernel_roofline_rows, + write_kernel_roofline_csv, +) + +KERNELS = [ + { + "name": "kernelA", + "pctRuntime": 60.0, + "count": 4.0, + "totalTime": 1000.0, + "points": [ + { + "peak": "HBM", + "ai": 2.0, + "perf": 500.0, + "roofPerf": 1000.0, + "pctRoof": 50.0, + "bandwidth": 500.0, + }, + { + "peak": "L2", + "ai": 4.0, + "perf": 500.0, + "roofPerf": None, + "pctRoof": None, + "bandwidth": None, + }, + ], + }, + { + "name": "kernelB", + "pctRuntime": 40.0, + "count": None, + "totalTime": None, + "points": [], + }, +] + + +def test_collect_kernel_roofline_rows_flattens_one_row_per_point() -> None: + """One row per kernel per plotted memory level; a kernel with no points + contributes no rows at all.""" + rows = collect_kernel_roofline_rows(KERNELS, time_unit="ns") + + assert [row["Kernel_Name"] for row in rows] == ["kernelA", "kernelA"] + assert [row["Cache_Level"] for row in rows] == ["HBM", "L2"] + assert rows[0]["Arithmetic_Intensity"] == 2.0 + assert rows[0]["Performance_GFLOPs"] == 500.0 + assert rows[0]["Roofline_Peak_GFLOPs"] == 1000.0 + assert rows[0]["Pct_of_Roofline"] == 50.0 + assert rows[0]["Bandwidth_GBps"] == 500.0 + assert rows[0]["Count"] == 4.0 + assert rows[0]["Total_Time"] == 1000.0 + assert rows[0]["Time_Unit"] == "ns" + assert rows[0]["Pct_Runtime"] == 60.0 + + # An unroofed point still gets a row, with the missing values as None. + assert rows[1]["Roofline_Peak_GFLOPs"] is None + assert rows[1]["Bandwidth_GBps"] is None + + +def test_collect_kernel_roofline_rows_empty_kernels_produce_no_rows() -> None: + assert collect_kernel_roofline_rows([], time_unit="ns") == [] + + +def test_build_kernel_roofline_dataframe_uses_fixed_column_order() -> None: + rows = collect_kernel_roofline_rows(KERNELS, time_unit="ns") + + df = build_kernel_roofline_dataframe(rows) + + assert list(df.columns) == KERNEL_ROOFLINE_CSV_COLUMNS + assert len(df) == 2 + + +def test_build_kernel_roofline_dataframe_empty_rows_keeps_columns() -> None: + """An empty rows list still yields a DataFrame with the expected columns, + so callers do not need to special-case an empty result.""" + df = build_kernel_roofline_dataframe([]) + + assert list(df.columns) == KERNEL_ROOFLINE_CSV_COLUMNS + assert len(df) == 0 + + +def test_write_kernel_roofline_csv_writes_expected_file(tmp_path: Path) -> None: + df = pd.DataFrame( + [{"Kernel_Name": "kernelA", "Cache_Level": "HBM"}], + columns=["Kernel_Name", "Cache_Level"], + ) + + output_path = write_kernel_roofline_csv(df, tmp_path) + + assert output_path == tmp_path / KERNEL_ROOFLINE_CSV_FILENAME + assert output_path.is_file() + written = pd.read_csv(output_path) + assert written.loc[0, "Kernel_Name"] == "kernelA" + assert written.loc[0, "Cache_Level"] == "HBM" diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index fea041b9305f..40693c74289b 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -13,10 +13,12 @@ from pathlib import Path from typing import TYPE_CHECKING, Dict, List, Tuple +import pandas as pd import plotly.graph_objects as go import pytest import roofline.roofline_html as roofline_html +from roofline.roofline_csv import KERNEL_ROOFLINE_CSV_FILENAME from roofline.roofline_frame import FRAME_X_MIN, canonical_frame from roofline.roofline_hover import wrap_hover_name from roofline.roofline_html import RooflineViewModel, build_interactive_document @@ -225,6 +227,42 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: assert unroofed[0]["points"][0]["hoverCells"] == ["N/A", "N/A"] +def test_kernel_points_carry_raw_values_matching_hover_cells() -> None: + """roofline_csv reads raw numeric point values directly (not the formatted + hover strings), so they have to be present and agree with what the + tooltip shows.""" + _, model = kernel_traces( + make_roofline(["FP32"]), + {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]}, + ) + point = model[0]["points"][0] + + assert point["roofPerf"] == pytest.approx( + float(point["hoverCells"][0].replace(",", "")) + ) + assert point["pctRoof"] == pytest.approx(float(point["hoverCells"][1])) + assert point["bandwidth"] == pytest.approx(CEILING["hbm"][2]) + + +def test_kernel_model_carries_count_and_total_time() -> None: + """The per-kernel view model keeps Count/Total_Time alongside pctRuntime, + since the CSV export reads all three straight off it.""" + _, model = kernel_traces( + make_roofline(["FP32"]), + { + "ai_hbm": [[100.0], [50000.0]], + "kernelNames": ["kA"], + "counts": [7.0], + "totalTime": [123.0], + "pctRuntime": [55.0], + }, + ) + + assert model[0]["count"] == 7.0 + assert model[0]["totalTime"] == 123.0 + assert model[0]["pctRuntime"] == 55.0 + + def test_kernel_hover_carries_the_whole_name() -> None: """A long demangled name reaches the tooltip whole. It is wrapped onto as many lines as it takes, but nothing is dropped: two instantiations of the From ddee08534870954d31bdf54813198600ea68f7c7 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Tue, 22 Sep 2026 20:28:57 +0100 Subject: [PATCH 07/45] bandwidth updated to show a value for a kernel --- .../rocprofiler-compute/src/roofline/roofline_main.py | 8 +++++--- .../tests/unit/roofline/test_roofline_main.py | 5 +++-- 2 files changed, 8 insertions(+), 5 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 793befd94d1d..b1ed56fb703f 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -498,13 +498,15 @@ def _build_kernel_points( continue cache_key = cache_level.removeprefix("ai_") - bandwidth = self._peak_value(ceiling_data, cache_key) roof_perf = self._roof_value_at( ai_value=ai_value, cache_key=cache_key, ceiling_data=ceiling_data, cap=compute_cap, ) + # This kernel's own achieved bandwidth at this level (not the + # hardware ceiling): performance (FLOP/s) / AI (FLOP/Byte) = Byte/s. + achieved_bandwidth = performance / ai_value pct_roof = 100.0 * performance / roof_perf if roof_perf else None points.append({ "peak": level_name, @@ -512,11 +514,11 @@ def _build_kernel_points( "perf": performance, "roofPerf": roof_perf, "pctRoof": pct_roof, - "bandwidth": bandwidth, + "bandwidth": achieved_bandwidth, "hoverCells": [ format_hover_number(roof_perf, ",.3f"), format_hover_number(pct_roof, ".4f"), - format_bandwidth(bandwidth) if bandwidth is not None else "N/A", + format_bandwidth(achieved_bandwidth), ], }) level_ai[level_name] = ai_value diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 40693c74289b..753151156fb3 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -224,7 +224,7 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: compute_peaks=[], ) assert "Performance limiter: Unknown" in unroofed_traces[0].hovertemplate - assert unroofed[0]["points"][0]["hoverCells"] == ["N/A", "N/A"] + assert unroofed[0]["points"][0]["hoverCells"] == ["N/A", "N/A", "900.000 GB/s"] def test_kernel_points_carry_raw_values_matching_hover_cells() -> None: @@ -241,7 +241,8 @@ def test_kernel_points_carry_raw_values_matching_hover_cells() -> None: float(point["hoverCells"][0].replace(",", "")) ) assert point["pctRoof"] == pytest.approx(float(point["hoverCells"][1])) - assert point["bandwidth"] == pytest.approx(CEILING["hbm"][2]) + # Achieved bandwidth for this kernel: performance / AI, not the ceiling. + assert point["bandwidth"] == pytest.approx(50000.0 / 100.0) def test_kernel_model_carries_count_and_total_time() -> None: From ed8a63f13cd40f93477285489c99597fc430f63c Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 10:45:52 +0100 Subject: [PATCH 08/45] Revert "roofline exported as csv" This reverts commit 2c3e908426c378a43996bf832e045b5e1f7d03a9. --- .../rocprof_compute_analyze/analysis_cli.py | 1 - .../src/roofline/roofline_csv.py | 69 ----------- .../src/roofline/roofline_main.py | 30 ----- .../integration/test_roofline_analyze.py | 33 ------ .../tests/unit/roofline/test_roofline_csv.py | 111 ------------------ .../tests/unit/roofline/test_roofline_main.py | 39 ------ 6 files changed, 283 deletions(-) delete mode 100644 projects/rocprofiler-compute/src/roofline/roofline_csv.py delete mode 100644 projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py diff --git a/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py b/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py index c657c61552d2..4d882c0ecb16 100644 --- a/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py +++ b/projects/rocprofiler-compute/src/rocprof_compute_analyze/analysis_cli.py @@ -244,7 +244,6 @@ def run_analysis(self) -> None: flops_dt, ) = roof_obj.construct_plotly_figures(ai_data=ai_data) roof_obj.save_html_files(ops_fig, flops_fig, ops_dt, flops_dt) - roof_obj.save_kernel_csv() else: console_warning( "roofline", diff --git a/projects/rocprofiler-compute/src/roofline/roofline_csv.py b/projects/rocprofiler-compute/src/roofline/roofline_csv.py deleted file mode 100644 index 3e6c52136b54..000000000000 --- a/projects/rocprofiler-compute/src/roofline/roofline_csv.py +++ /dev/null @@ -1,69 +0,0 @@ -# Copyright (c) Advanced Micro Devices, Inc. -# SPDX-License-Identifier: MIT - -"""Per-kernel roofline data export, alongside the interactive roofline HTML.""" - -from pathlib import Path -from typing import Any, Union - -import pandas as pd - -from utils.logger import console_log - -KERNEL_ROOFLINE_CSV_FILENAME = "roofline_kernels.csv" - -KERNEL_ROOFLINE_CSV_COLUMNS = [ - "Kernel_Name", - "Cache_Level", - "Arithmetic_Intensity", - "Performance_GFLOPs", - "Roofline_Peak_GFLOPs", - "Pct_of_Roofline", - "Bandwidth_GBps", - "Count", - "Total_Time", - "Time_Unit", - "Pct_Runtime", -] - - -def collect_kernel_roofline_rows( - kernels: list[dict[str, Any]], time_unit: str -) -> list[dict[str, Any]]: - """Flatten the FLOP figure's per-kernel view-model data into tidy rows. - - One row per kernel per memory level actually plotted, read straight from - RooflineViewModel.kernels (built by Roofline._build_kernel_traces), so - these numbers can never diverge from what the interactive HTML shows. - """ - rows: list[dict[str, Any]] = [] - for kernel in kernels: - for point in kernel["points"]: - rows.append({ - "Kernel_Name": kernel["name"], - "Cache_Level": point["peak"], - "Arithmetic_Intensity": point["ai"], - "Performance_GFLOPs": point["perf"], - "Roofline_Peak_GFLOPs": point["roofPerf"], - "Pct_of_Roofline": point["pctRoof"], - "Bandwidth_GBps": point["bandwidth"], - "Count": kernel.get("count"), - "Total_Time": kernel.get("totalTime"), - "Time_Unit": time_unit, - "Pct_Runtime": kernel.get("pctRuntime"), - }) - return rows - - -def build_kernel_roofline_dataframe(rows: list[dict[str, Any]]) -> pd.DataFrame: - """Assemble tidy per-kernel roofline rows into a DataFrame with a fixed - column order.""" - return pd.DataFrame(rows, columns=KERNEL_ROOFLINE_CSV_COLUMNS) - - -def write_kernel_roofline_csv(df: pd.DataFrame, workload_dir: Union[str, Path]) -> Path: - """Write the per-kernel roofline DataFrame to roofline_kernels.csv.""" - output_file = Path(workload_dir) / KERNEL_ROOFLINE_CSV_FILENAME - df.to_csv(output_file, index=False) - console_log("roofline", f"Saved per-kernel roofline data to {output_file}") - return output_file diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index b1ed56fb703f..d4a74099c823 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -12,11 +12,6 @@ import plotly.colors as pcolors import plotly.graph_objects as go -from roofline.roofline_csv import ( - build_kernel_roofline_dataframe, - collect_kernel_roofline_rows, - write_kernel_roofline_csv, -) from roofline.roofline_frame import canonical_frame from roofline.roofline_hover import ( build_compute_peak_hover, @@ -463,8 +458,6 @@ def _build_kernel_traces( "color": color, "points": points, "pctRuntime": pct_val, - "count": count_val, - "totalTime": time_val, }) return traces, kernels_model @@ -512,9 +505,6 @@ def _build_kernel_points( "peak": level_name, "ai": ai_value, "perf": performance, - "roofPerf": roof_perf, - "pctRoof": pct_roof, - "bandwidth": achieved_bandwidth, "hoverCells": [ format_hover_number(roof_perf, ",.3f"), format_hover_number(pct_roof, ".4f"), @@ -625,26 +615,6 @@ def save_html_files( Path(path).write_text(document, encoding="utf-8") console_log("roofline", "Roofline HTML files saved.") - def save_kernel_csv(self) -> Optional[Path]: - """Write the per-kernel roofline data next to the HTML, as a tidy CSV. - - Reuses the FLOP figure's kernel view-model, so the CSV and the - interactive HTML always agree. Call after construct_plotly_figures, - which is what populates that view-model. - """ - view_model = self.__view_models.get("FLOP") - if view_model is None or not view_model.kernels: - return None - - time_unit = (self.__ai_data or {}).get("timeUnit", "") - rows = collect_kernel_roofline_rows(view_model.kernels, time_unit) - if not rows: - return None - - df = build_kernel_roofline_dataframe(rows) - workload_dir = self.__run_parameters["workload_dir"] - return write_kernel_roofline_csv(df, workload_dir) - def _combined_html_figure( self, ops_figure: Optional[go.Figure], diff --git a/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py b/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py index 3c19451a8bc0..137624654bbd 100644 --- a/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py +++ b/projects/rocprofiler-compute/tests/integration/test_roofline_analyze.py @@ -17,7 +17,6 @@ import common import pytest -from roofline.roofline_csv import KERNEL_ROOFLINE_CSV_FILENAME from tests.integration import common as integration_common config = {} @@ -75,38 +74,6 @@ def test_analyze_generates_roofline_html( common.clean_output_dir(config["cleanup"], workload_dir) -def test_analyze_generates_kernel_roofline_csv( - binary_handler_analyze_rocprof_compute: Callable[[list[str]], int], -) -> None: - """ - Analyze writes roofline_kernels.csv alongside the roofline HTML, with the - same per-kernel data plotted in the interactive figure. - - The MI200 fixture's kernel is a pure memory-bandwidth microbenchmark with - zero measured FLOPs, so it draws no dot on the FLOP figure and no CSV row - either -- matching the empty-HTML-plot case exactly, and confirming the - export doesn't write a spurious empty file. Coverage for the populated - (non-empty rows, expected columns) case lives in the unit tests, where - kernel performance data can be constructed directly. - """ - workload_dir = integration_common.setup_workload_dir(roofline_dir) - - assert (Path(workload_dir) / "roofline.csv").exists() - - code = binary_handler_analyze_rocprof_compute([ - "analyze", - "--path", - workload_dir, - "--roofline-data-type", - "FP32", - ]) - assert code == 0 - - assert not (Path(workload_dir) / KERNEL_ROOFLINE_CSV_FILENAME).exists() - - common.clean_output_dir(config["cleanup"], workload_dir) - - def test_analyze_roofline_datatype_independently( binary_handler_analyze_rocprof_compute: Callable[[list[str]], int], ) -> None: diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py deleted file mode 100644 index b60099abaf01..000000000000 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_csv.py +++ /dev/null @@ -1,111 +0,0 @@ -# Copyright (c) Advanced Micro Devices, Inc. -# SPDX-License-Identifier: MIT - -"""Unit coverage for the per-kernel roofline CSV export: flattening the FLOP -figure's kernel view-model into tidy rows and writing them to disk. -""" - -from pathlib import Path - -import pandas as pd - -from roofline.roofline_csv import ( - KERNEL_ROOFLINE_CSV_COLUMNS, - KERNEL_ROOFLINE_CSV_FILENAME, - build_kernel_roofline_dataframe, - collect_kernel_roofline_rows, - write_kernel_roofline_csv, -) - -KERNELS = [ - { - "name": "kernelA", - "pctRuntime": 60.0, - "count": 4.0, - "totalTime": 1000.0, - "points": [ - { - "peak": "HBM", - "ai": 2.0, - "perf": 500.0, - "roofPerf": 1000.0, - "pctRoof": 50.0, - "bandwidth": 500.0, - }, - { - "peak": "L2", - "ai": 4.0, - "perf": 500.0, - "roofPerf": None, - "pctRoof": None, - "bandwidth": None, - }, - ], - }, - { - "name": "kernelB", - "pctRuntime": 40.0, - "count": None, - "totalTime": None, - "points": [], - }, -] - - -def test_collect_kernel_roofline_rows_flattens_one_row_per_point() -> None: - """One row per kernel per plotted memory level; a kernel with no points - contributes no rows at all.""" - rows = collect_kernel_roofline_rows(KERNELS, time_unit="ns") - - assert [row["Kernel_Name"] for row in rows] == ["kernelA", "kernelA"] - assert [row["Cache_Level"] for row in rows] == ["HBM", "L2"] - assert rows[0]["Arithmetic_Intensity"] == 2.0 - assert rows[0]["Performance_GFLOPs"] == 500.0 - assert rows[0]["Roofline_Peak_GFLOPs"] == 1000.0 - assert rows[0]["Pct_of_Roofline"] == 50.0 - assert rows[0]["Bandwidth_GBps"] == 500.0 - assert rows[0]["Count"] == 4.0 - assert rows[0]["Total_Time"] == 1000.0 - assert rows[0]["Time_Unit"] == "ns" - assert rows[0]["Pct_Runtime"] == 60.0 - - # An unroofed point still gets a row, with the missing values as None. - assert rows[1]["Roofline_Peak_GFLOPs"] is None - assert rows[1]["Bandwidth_GBps"] is None - - -def test_collect_kernel_roofline_rows_empty_kernels_produce_no_rows() -> None: - assert collect_kernel_roofline_rows([], time_unit="ns") == [] - - -def test_build_kernel_roofline_dataframe_uses_fixed_column_order() -> None: - rows = collect_kernel_roofline_rows(KERNELS, time_unit="ns") - - df = build_kernel_roofline_dataframe(rows) - - assert list(df.columns) == KERNEL_ROOFLINE_CSV_COLUMNS - assert len(df) == 2 - - -def test_build_kernel_roofline_dataframe_empty_rows_keeps_columns() -> None: - """An empty rows list still yields a DataFrame with the expected columns, - so callers do not need to special-case an empty result.""" - df = build_kernel_roofline_dataframe([]) - - assert list(df.columns) == KERNEL_ROOFLINE_CSV_COLUMNS - assert len(df) == 0 - - -def test_write_kernel_roofline_csv_writes_expected_file(tmp_path: Path) -> None: - df = pd.DataFrame( - [{"Kernel_Name": "kernelA", "Cache_Level": "HBM"}], - columns=["Kernel_Name", "Cache_Level"], - ) - - output_path = write_kernel_roofline_csv(df, tmp_path) - - assert output_path == tmp_path / KERNEL_ROOFLINE_CSV_FILENAME - assert output_path.is_file() - written = pd.read_csv(output_path) - assert written.loc[0, "Kernel_Name"] == "kernelA" - assert written.loc[0, "Cache_Level"] == "HBM" diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 753151156fb3..fefa1a744fd7 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -13,12 +13,10 @@ from pathlib import Path from typing import TYPE_CHECKING, Dict, List, Tuple -import pandas as pd import plotly.graph_objects as go import pytest import roofline.roofline_html as roofline_html -from roofline.roofline_csv import KERNEL_ROOFLINE_CSV_FILENAME from roofline.roofline_frame import FRAME_X_MIN, canonical_frame from roofline.roofline_hover import wrap_hover_name from roofline.roofline_html import RooflineViewModel, build_interactive_document @@ -227,43 +225,6 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: assert unroofed[0]["points"][0]["hoverCells"] == ["N/A", "N/A", "900.000 GB/s"] -def test_kernel_points_carry_raw_values_matching_hover_cells() -> None: - """roofline_csv reads raw numeric point values directly (not the formatted - hover strings), so they have to be present and agree with what the - tooltip shows.""" - _, model = kernel_traces( - make_roofline(["FP32"]), - {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]}, - ) - point = model[0]["points"][0] - - assert point["roofPerf"] == pytest.approx( - float(point["hoverCells"][0].replace(",", "")) - ) - assert point["pctRoof"] == pytest.approx(float(point["hoverCells"][1])) - # Achieved bandwidth for this kernel: performance / AI, not the ceiling. - assert point["bandwidth"] == pytest.approx(50000.0 / 100.0) - - -def test_kernel_model_carries_count_and_total_time() -> None: - """The per-kernel view model keeps Count/Total_Time alongside pctRuntime, - since the CSV export reads all three straight off it.""" - _, model = kernel_traces( - make_roofline(["FP32"]), - { - "ai_hbm": [[100.0], [50000.0]], - "kernelNames": ["kA"], - "counts": [7.0], - "totalTime": [123.0], - "pctRuntime": [55.0], - }, - ) - - assert model[0]["count"] == 7.0 - assert model[0]["totalTime"] == 123.0 - assert model[0]["pctRuntime"] == 55.0 - - def test_kernel_hover_carries_the_whole_name() -> None: """A long demangled name reaches the tooltip whole. It is wrapped onto as many lines as it takes, but nothing is dropped: two instantiations of the From 32c86484936ea0b7eca6d7d6bb5b15aae4c7016e Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 11:31:14 +0100 Subject: [PATCH 09/45] kernel name truncation --- .../src/roofline/assets/roofline_plot.js | 14 ++++- .../src/roofline/roofline_hover.py | 13 ++++- .../src/roofline/roofline_main.py | 23 ++++---- .../tests/unit/roofline/test_roofline_main.py | 53 +++++++++++++++---- 4 files changed, 81 insertions(+), 22 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 889c7ba86df1..97248cbb287d 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -22,6 +22,7 @@ // ---- Own presentation --------------------------------------------------- var ALL_PEAKS_LABEL = "All peaks"; + var KERNEL_LIST_NAME_MAX_LENGTH = 200; var FALLBACK_COLOR = "#888888"; var PLOT_DIM_OPACITY = 0.15; // How far a roof drawn at 45 degrees in data space may lean on screen. @@ -205,6 +206,13 @@ return "(" + shown + " / " + total + ")"; } + function truncateKernelListName(name) { + if (name.length <= KERNEL_LIST_NAME_MAX_LENGTH) { + return name; + } + return name.slice(0, KERNEL_LIST_NAME_MAX_LENGTH - 3) + "..."; + } + function eachKernelRow(fn) { if (!kernelList) { return; @@ -1408,6 +1416,9 @@ var label = document.createElement("span"); label.className = opts.labelClass || "roofline-panel-name"; label.textContent = opts.label; + if (opts.title) { + label.title = opts.title; + } action.appendChild(swatch); action.appendChild(label); @@ -1508,7 +1519,8 @@ kernelList.appendChild( createPanelRow({ color: kernel.color, - label: kernel.name, + label: truncateKernelListName(kernel.name), + title: kernel.name, labelClass: "roofline-panel-name roofline-kernel-name", dataset: { index: String(index) }, actionExtras: actionExtras, diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index 53805ac30b47..258d27afd928 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -11,6 +11,16 @@ _HOVER_WRAP_WIDTH = 44 +def truncate_kernel_name(name: str) -> str: + """Cap a kernel name at 200 characters so a pathological name can't blow + up the tooltip.""" + max_length = 200 + suffix = "..." + if len(name) <= max_length: + return name + return name[: max_length - len(suffix)] + suffix + + def wrap_hover_name(name: str) -> str: """Wrap a kernel name across as many tooltip lines as it takes.""" if not name: @@ -29,6 +39,7 @@ def wrap_hover_name(name: str) -> str: def build_kernel_hover_template( name_html: str, limiter: str, + limiter_category: str, count: Optional[float], total_time: Optional[float], time_unit: str, @@ -45,12 +56,12 @@ def build_kernel_hover_template( return _hover( name_html, [ + f"Limited by {limiter_category}: {limiter}", "AI: %{x:.6g}", f"Achieved throughput: %{{y:,.3f}} {unit}", f"Peak throughput: %{{customdata[0]}} {unit}", "Percent of roofline achieved: %{customdata[1]} %", "Cache level bandwidth: %{customdata[2]}", - f"Performance limiter: {limiter}", f"Total dispatches: {_format_integer(count)}", f"Aggregate time in kernel: {time_txt}", f"Aggregate percent runtime: {format_hover_number(pct_runtime, '.5f')} %", diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index d4a74099c823..68d5f2c88dbd 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -19,6 +19,7 @@ build_roof_hover, format_bandwidth, format_hover_number, + truncate_kernel_name, wrap_hover_name, ) from roofline.roofline_html import ( @@ -373,24 +374,27 @@ def _determine_kernel_limiter( ceiling_data: dict[str, Any], compute_cap: float, compute_cap_label: str, - ) -> str: + ) -> tuple[str, str]: """Name the specific binding roof for a kernel: the roof with the lowest achievable performance at the kernel's operating point. The compute candidate is the envelope cap the diagonals are actually drawn to, so the limiter agrees with the drawn roof and with the percent of roofline - the tooltip reports.""" - candidates: list[tuple[float, str]] = [] + the tooltip reports. + + Returns (label, category), where category is "Memory" or "Compute".""" + candidates: list[tuple[float, str, str]] = [] for level_name, ai_value in level_ai.items(): bandwidth = self._peak_value(ceiling_data, level_name.lower()) if bandwidth and ai_value > 0: - candidates.append((bandwidth * ai_value, level_name)) + candidates.append((bandwidth * ai_value, level_name, "Memory")) if compute_cap != float("inf"): - candidates.append((compute_cap, compute_cap_label)) + candidates.append((compute_cap, compute_cap_label, "Compute")) if not candidates: - return "Unknown" - return min(candidates, key=lambda candidate: candidate[0])[1] + return "Unknown", "Unknown" + _, label, category = min(candidates, key=lambda candidate: candidate[0]) + return label, category def _build_kernel_traces( self, @@ -425,7 +429,7 @@ def _build_kernel_traces( values[kernel_index] if kernel_index < len(values) else None for values in (kernel_colors, counts, total_time, pct_runtime) ) - limiter = self._determine_kernel_limiter( + limiter, limiter_category = self._determine_kernel_limiter( level_ai, ceiling_data, compute_cap, compute_cap_label ) @@ -443,8 +447,9 @@ def _build_kernel_traces( ), customdata=[point["hoverCells"] for point in points], hovertemplate=build_kernel_hover_template( - name_html=wrap_hover_name(kernel_name), + name_html=wrap_hover_name(truncate_kernel_name(kernel_name)), limiter=limiter, + limiter_category=limiter_category, count=count_val, total_time=time_val, time_unit=time_unit, diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index fefa1a744fd7..c51e33d94a3e 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -18,7 +18,7 @@ import roofline.roofline_html as roofline_html from roofline.roofline_frame import FRAME_X_MIN, canonical_frame -from roofline.roofline_hover import wrap_hover_name +from roofline.roofline_hover import truncate_kernel_name, wrap_hover_name from roofline.roofline_html import RooflineViewModel, build_interactive_document if TYPE_CHECKING: @@ -194,8 +194,8 @@ def test_kernel_traces_score_against_the_tallest_drawn_ceiling() -> None: assert pct_roof(matrix_capped[0]) < 100.0 assert pct_roof(valu_capped[0]) > 100.0 - assert "Performance limiter: FP32 MFMA" in matrix_traces[0].hovertemplate - assert "Performance limiter: FP32 VALU" in valu_traces[0].hovertemplate + assert "Limited by Compute: FP32 MFMA" in matrix_traces[0].hovertemplate + assert "Limited by Compute: FP32 VALU" in valu_traces[0].hovertemplate def test_kernel_traces_name_the_roof_that_binds() -> None: @@ -213,7 +213,7 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: ) assert [kernel["name"] for kernel in model] == ["kA"] assert [point["peak"] for point in model[0]["points"]] == ["HBM"] - assert "Performance limiter: HBM" in traces[0].hovertemplate + assert "Limited by Memory: HBM" in traces[0].hovertemplate unroofed_traces, unroofed = kernel_traces( make_roofline(["FP32"]), @@ -221,16 +221,16 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: ceiling_data={}, compute_peaks=[], ) - assert "Performance limiter: Unknown" in unroofed_traces[0].hovertemplate + assert "Limited by Unknown: Unknown" in unroofed_traces[0].hovertemplate assert unroofed[0]["points"][0]["hoverCells"] == ["N/A", "N/A", "900.000 GB/s"] -def test_kernel_hover_carries_the_whole_name() -> None: - """A long demangled name reaches the tooltip whole. It is wrapped onto as - many lines as it takes, but nothing is dropped: two instantiations of the - same function are told apart by template arguments that run to the very end - of the name.""" - name = "Cijk_Alik_Bljk_" + "SB_MT256x256x16_MI32x32x2x1_" * 40 +def test_kernel_hover_carries_the_whole_name_up_to_the_limit() -> None: + """A demangled name within the tooltip's length cap reaches the tooltip + whole. It is wrapped onto as many lines as it takes, but nothing is + dropped: two instantiations of the same function are told apart by + template arguments that run to the very end of the name.""" + name = "Cijk_Alik_Bljk_" + "SB_MT256x256x16_MI32x32x2x1_" * 6 traces, _ = kernel_traces( make_roofline(["FP32"]), @@ -246,6 +246,37 @@ def test_kernel_hover_carries_the_whole_name() -> None: assert lines.replace("
", "") == name +def test_truncate_kernel_name_caps_pathologically_long_names() -> None: + """A pathologically long demangled name is capped at 200 characters, ending + in an ellipsis, so it can't blow up the tooltip. Names within the cap are + untouched.""" + name = "Cijk_Alik_Bljk_" + "SB_MT256x256x16_MI32x32x2x1_" * 40 + + truncated = truncate_kernel_name(name) + + assert len(truncated) == 200 + assert truncated == name[:197] + "..." + assert truncate_kernel_name("short_name") == "short_name" + + +def test_kernel_hover_truncates_names_beyond_two_hundred_characters() -> None: + """The kernel trace hover wraps the truncated name, not the raw one, so a + pathologically long demangled name can't blow up the tooltip.""" + name = "Cijk_Alik_Bljk_" + "SB_MT256x256x16_MI32x32x2x1_" * 40 + + traces, _ = kernel_traces( + make_roofline(["FP32"]), + {"ai_hbm": [[1.0], [900.0]], "kernelNames": [name]}, + ) + + wrapped = wrap_hover_name(truncate_kernel_name(name)) + assert wrapped in traces[0].hovertemplate + lines = wrapped.split(">", 1)[1].removesuffix("") + flattened = lines.replace("
", "") + assert len(flattened) == 200 + assert flattened.endswith("...") + + def test_canonical_frame_bounds_are_whole_decades() -> None: """Every axis limit lands on a power of ten.""" bounds = canonical_frame([5300.0, 10000.0], [81000.0, 163000.0]) From 26b387b11ed07bfa1e7c9f8618f59502f6071834 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 12:05:44 +0100 Subject: [PATCH 10/45] kernel name hover for kernel list --- .../src/roofline/assets/roofline_plot.css | 28 +++++++ .../src/roofline/assets/roofline_plot.js | 78 ++++++++++++++++++- 2 files changed, 104 insertions(+), 2 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css index 583becfa808e..0aa6105afba0 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css @@ -499,6 +499,34 @@ body { cursor: pointer; } +.roofline-name-tooltip { + position: fixed; + top: 0; + left: 0; + z-index: 500; + max-width: min(480px, calc(100vw - 24px)); + padding: 6px 10px; + background: var(--roofline-surface); + color: var(--roofline-text); + border: 1px solid var(--roofline-border); + border-radius: 6px; + box-shadow: 0 4px 12px rgba(0, 0, 0, 0.15); + font-family: var(--roofline-kernel-font, monospace); + font-size: 12px; + line-height: 1.35; + word-break: break-word; + overflow-wrap: anywhere; + pointer-events: none; + opacity: 0; + visibility: hidden; + transition: opacity 0.1s ease; +} + +.roofline-name-tooltip.visible { + opacity: 1; + visibility: visible; +} + @media (max-width: 900px) { .roofline-body { flex-direction: column; diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 97248cbb287d..674f61c059e3 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -213,6 +213,57 @@ return name.slice(0, KERNEL_LIST_NAME_MAX_LENGTH - 3) + "..."; } + // A native `title` attribute gets dismissed by the row's click-to-filter + // handler and won't reappear until the pointer leaves and re-enters, so the + // full kernel name uses this mouse-following tooltip instead. + var nameTooltip = null; + var nameTooltipTimer = null; + var NAME_TOOLTIP_OFFSET = 14; + var NAME_TOOLTIP_DELAY_MS = 2000; + + function ensureNameTooltip() { + if (!nameTooltip) { + nameTooltip = document.createElement("div"); + nameTooltip.className = "roofline-name-tooltip"; + nameTooltip.setAttribute("role", "tooltip"); + document.body.appendChild(nameTooltip); + } + return nameTooltip; + } + + function positionNameTooltip(x, y) { + var tooltip = nameTooltip; + if (!tooltip) { + return; + } + var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); + var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); + tooltip.style.left = clamp(x + NAME_TOOLTIP_OFFSET, 4, maxLeft) + "px"; + tooltip.style.top = clamp(y + NAME_TOOLTIP_OFFSET, 4, maxTop) + "px"; + } + + function showNameTooltip(text, x, y) { + var tooltip = ensureNameTooltip(); + tooltip.textContent = text; + tooltip.classList.add("visible"); + positionNameTooltip(x, y); + } + + function scheduleNameTooltip(text, getPosition) { + clearTimeout(nameTooltipTimer); + nameTooltipTimer = setTimeout(function () { + var position = getPosition(); + showNameTooltip(text, position.x, position.y); + }, NAME_TOOLTIP_DELAY_MS); + } + + function hideNameTooltip() { + clearTimeout(nameTooltipTimer); + if (nameTooltip) { + nameTooltip.classList.remove("visible"); + } + } + function eachKernelRow(fn) { if (!kernelList) { return; @@ -1417,7 +1468,27 @@ label.className = opts.labelClass || "roofline-panel-name"; label.textContent = opts.label; if (opts.title) { - label.title = opts.title; + var lastMousePosition = { x: 0, y: 0 }; + label.addEventListener("mouseenter", function (event) { + lastMousePosition = { x: event.clientX, y: event.clientY }; + scheduleNameTooltip(opts.title, function () { + return lastMousePosition; + }); + }); + label.addEventListener("mousemove", function (event) { + lastMousePosition = { x: event.clientX, y: event.clientY }; + if (nameTooltip && nameTooltip.classList.contains("visible")) { + positionNameTooltip(event.clientX, event.clientY); + } + }); + label.addEventListener("mouseleave", hideNameTooltip); + action.addEventListener("focus", function () { + var rect = label.getBoundingClientRect(); + scheduleNameTooltip(opts.title, function () { + return { x: rect.left, y: rect.bottom }; + }); + }); + action.addEventListener("blur", hideNameTooltip); } action.appendChild(swatch); @@ -1425,7 +1496,10 @@ opts.actionExtras.forEach(function (node) { action.appendChild(node); }); - action.addEventListener("click", opts.onClick); + action.addEventListener("click", function (event) { + hideNameTooltip(); + opts.onClick(event); + }); item.appendChild(action); opts.siblingControls.forEach(function (node) { item.appendChild(node); From d909ac743273952645732add78fafec1832d2e52 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 18:20:01 +0100 Subject: [PATCH 11/45] Hover data refactoring --- projects/rocprofiler-compute/src/roofline/roofline_hover.py | 4 +--- projects/rocprofiler-compute/src/roofline/roofline_main.py | 4 ++-- 2 files changed, 3 insertions(+), 5 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index 258d27afd928..b00cad6afddb 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -57,10 +57,8 @@ def build_kernel_hover_template( name_html, [ f"Limited by {limiter_category}: {limiter}", + f"Performance: %{{customdata[1]}}% (%{{y:,.0f}} / %{{customdata[0]}} {unit})", "AI: %{x:.6g}", - f"Achieved throughput: %{{y:,.3f}} {unit}", - f"Peak throughput: %{{customdata[0]}} {unit}", - "Percent of roofline achieved: %{customdata[1]} %", "Cache level bandwidth: %{customdata[2]}", f"Total dispatches: {_format_integer(count)}", f"Aggregate time in kernel: {time_txt}", diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 68d5f2c88dbd..d29c41049149 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -511,8 +511,8 @@ def _build_kernel_points( "ai": ai_value, "perf": performance, "hoverCells": [ - format_hover_number(roof_perf, ",.3f"), - format_hover_number(pct_roof, ".4f"), + format_hover_number(roof_perf, ",.0f"), + format_hover_number(pct_roof, ",.2f"), format_bandwidth(achieved_bandwidth), ], }) From 2ab6904e614ad78d0822ab5aaeaa00c0cbb52ef1 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 18:35:25 +0100 Subject: [PATCH 12/45] Hover updates --- .../src/roofline/roofline_hover.py | 25 ++++++---- .../src/roofline/roofline_main.py | 48 +++++++++++++++++-- .../tests/unit/roofline/test_roofline_main.py | 6 ++- 3 files changed, 65 insertions(+), 14 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index b00cad6afddb..4d44f849d01f 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -41,6 +41,7 @@ def build_kernel_hover_template( limiter: str, limiter_category: str, count: Optional[float], + total_dispatches: Optional[float], total_time: Optional[float], time_unit: str, pct_runtime: Optional[float], @@ -48,21 +49,29 @@ def build_kernel_hover_template( ) -> str: """Kernel hover template; per-point values come from customdata.""" unit = f"G{ops_flops}s/s" - time_txt = ( - f"{format_hover_number(total_time, ',.2f')} {time_unit}".strip() - if total_time is not None - else "N/A" + kernel_time_txt = format_hover_number(total_time, ",.0f") + total_app_time = ( + total_time / (pct_runtime / 100.0) + if total_time is not None and pct_runtime + else None ) + total_app_time_txt = format_hover_number(total_app_time, ",.0f") + pct_runtime_txt = format_hover_number(pct_runtime, ",.2f") + pct_dispatches = ( + 100.0 * count / total_dispatches + if count is not None and total_dispatches + else None + ) + pct_dispatches_txt = format_hover_number(pct_dispatches, ",.2f") return _hover( name_html, [ f"Limited by {limiter_category}: {limiter}", f"Performance: %{{customdata[1]}}% (%{{y:,.0f}} / %{{customdata[0]}} {unit})", + "%{customdata[2]}", + f"Dispatch Number: {pct_dispatches_txt}% ({_format_integer(count)} / {_format_integer(total_dispatches)})", + f"Duration: {pct_runtime_txt}% ({kernel_time_txt} / {total_app_time_txt} {time_unit})", "AI: %{x:.6g}", - "Cache level bandwidth: %{customdata[2]}", - f"Total dispatches: {_format_integer(count)}", - f"Aggregate time in kernel: {time_txt}", - f"Aggregate percent runtime: {format_hover_number(pct_runtime, '.5f')} %", ], ) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index d29c41049149..9d202ca7df7c 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -413,6 +413,7 @@ def _build_kernel_traces( total_time = self.__ai_data.get("totalTime", []) pct_runtime = self.__ai_data.get("pctRuntime", []) time_unit = self.__ai_data.get("timeUnit", "") + total_dispatches = sum(count for count in counts if count is not None) compute_cap, compute_cap_label = self._envelope_compute_cap(compute_peaks) for kernel_index, kernel_name in enumerate(kernel_names): @@ -451,6 +452,7 @@ def _build_kernel_traces( limiter=limiter, limiter_category=limiter_category, count=count_val, + total_dispatches=total_dispatches, total_time=time_val, time_unit=time_unit, pct_runtime=pct_val, @@ -505,21 +507,57 @@ def _build_kernel_points( # This kernel's own achieved bandwidth at this level (not the # hardware ceiling): performance (FLOP/s) / AI (FLOP/Byte) = Byte/s. achieved_bandwidth = performance / ai_value + peak_bandwidth = self._peak_value(ceiling_data, cache_key) pct_roof = 100.0 * performance / roof_perf if roof_perf else None points.append({ "peak": level_name, "ai": ai_value, "perf": performance, - "hoverCells": [ - format_hover_number(roof_perf, ",.0f"), - format_hover_number(pct_roof, ",.2f"), - format_bandwidth(achieved_bandwidth), - ], + "level_name": level_name, + "roof_perf_txt": format_hover_number(roof_perf, ",.0f"), + "pct_roof_txt": format_hover_number(pct_roof, ",.2f"), + "achieved_bandwidth": achieved_bandwidth, + "peak_bandwidth": peak_bandwidth, }) level_ai[level_name] = ai_value + bandwidth_html = self._build_bandwidth_hover_html(points) + for point in points: + point["hoverCells"] = [ + point.pop("roof_perf_txt"), + point.pop("pct_roof_txt"), + bandwidth_html, + ] + point.pop("level_name") + point.pop("achieved_bandwidth") + point.pop("peak_bandwidth") + return points, level_ai + @staticmethod + def _build_bandwidth_hover_html(points: list[dict[str, Any]]) -> str: + """One 'Bandwidth:' block listing every cache level this kernel has a + point for, each formatted like the Performance field: percent + (achieved/peak).""" + if not points: + return "Bandwidth: N/A" + lines = ["Bandwidth:"] + for point in points: + peak_bandwidth = point["peak_bandwidth"] + achieved_bandwidth = point["achieved_bandwidth"] + pct_bandwidth = ( + 100.0 * achieved_bandwidth / peak_bandwidth + if peak_bandwidth + else None + ) + lines.append( + f" {point['level_name']}: " + f"{format_hover_number(pct_bandwidth, ',.2f')}% " + f"({format_bandwidth(achieved_bandwidth)}/" + f"{format_bandwidth(peak_bandwidth)})" + ) + return "
".join(lines) + @demarcate def construct_plotly_figures( self, ai_data: dict[str, Any], datatypes: Optional[list[str]] = None diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index c51e33d94a3e..13247a79c226 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -222,7 +222,11 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: compute_peaks=[], ) assert "Limited by Unknown: Unknown" in unroofed_traces[0].hovertemplate - assert unroofed[0]["points"][0]["hoverCells"] == ["N/A", "N/A", "900.000 GB/s"] + assert unroofed[0]["points"][0]["hoverCells"] == [ + "N/A", + "N/A", + "Bandwidth:
 HBM: N/A% (900.000 GB/s/N/A)", + ] def test_kernel_hover_carries_the_whole_name_up_to_the_limit() -> None: From ab0a76175e10e7ae33dd9b3539b6079564a7158a Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 18:58:06 +0100 Subject: [PATCH 13/45] Show cache levels even if zero --- .../src/roofline/roofline_main.py | 41 +++++++++++-------- .../tests/unit/roofline/test_roofline_main.py | 25 ++++++++++- 2 files changed, 49 insertions(+), 17 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 9d202ca7df7c..a35f066d671a 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -481,6 +481,7 @@ def _build_kernel_points( Also returns the level -> AI map the limiter is chosen from. """ points: list[dict[str, Any]] = [] + bandwidth_entries: list[dict[str, Any]] = [] level_ai: dict[str, float] = {} for cache_level in CACHE_LEVELS: @@ -494,10 +495,19 @@ def _build_kernel_points( continue ai_value = level_points[0][kernel_index] performance = level_points[1][kernel_index] + cache_key = cache_level.removeprefix("ai_") + peak_bandwidth = self._peak_value(ceiling_data, cache_key) + if not (ai_value > 0 and performance > 0): + # A level with no traffic still gets a 0-valued bandwidth + # line instead of silently disappearing from the tooltip. + bandwidth_entries.append({ + "level_name": level_name, + "achieved_bandwidth": 0.0, + "peak_bandwidth": peak_bandwidth, + }) continue - cache_key = cache_level.removeprefix("ai_") roof_perf = self._roof_value_at( ai_value=ai_value, cache_key=cache_key, @@ -507,51 +517,50 @@ def _build_kernel_points( # This kernel's own achieved bandwidth at this level (not the # hardware ceiling): performance (FLOP/s) / AI (FLOP/Byte) = Byte/s. achieved_bandwidth = performance / ai_value - peak_bandwidth = self._peak_value(ceiling_data, cache_key) pct_roof = 100.0 * performance / roof_perf if roof_perf else None points.append({ "peak": level_name, "ai": ai_value, "perf": performance, - "level_name": level_name, "roof_perf_txt": format_hover_number(roof_perf, ",.0f"), "pct_roof_txt": format_hover_number(pct_roof, ",.2f"), + }) + bandwidth_entries.append({ + "level_name": level_name, "achieved_bandwidth": achieved_bandwidth, "peak_bandwidth": peak_bandwidth, }) level_ai[level_name] = ai_value - bandwidth_html = self._build_bandwidth_hover_html(points) + bandwidth_html = self._build_bandwidth_hover_html(bandwidth_entries) for point in points: point["hoverCells"] = [ point.pop("roof_perf_txt"), point.pop("pct_roof_txt"), bandwidth_html, ] - point.pop("level_name") - point.pop("achieved_bandwidth") - point.pop("peak_bandwidth") return points, level_ai @staticmethod - def _build_bandwidth_hover_html(points: list[dict[str, Any]]) -> str: - """One 'Bandwidth:' block listing every cache level this kernel has a - point for, each formatted like the Performance field: percent - (achieved/peak).""" - if not points: + def _build_bandwidth_hover_html(bandwidth_entries: list[dict[str, Any]]) -> str: + """One 'Bandwidth:' block listing every cache level this kernel has + data for, each formatted like the Performance field: percent + (achieved/peak). A level with no traffic still shows up at 0, + rather than vanishing from the tooltip.""" + if not bandwidth_entries: return "Bandwidth: N/A" lines = ["Bandwidth:"] - for point in points: - peak_bandwidth = point["peak_bandwidth"] - achieved_bandwidth = point["achieved_bandwidth"] + for entry in bandwidth_entries: + peak_bandwidth = entry["peak_bandwidth"] + achieved_bandwidth = entry["achieved_bandwidth"] pct_bandwidth = ( 100.0 * achieved_bandwidth / peak_bandwidth if peak_bandwidth else None ) lines.append( - f" {point['level_name']}: " + f"\u2003{entry['level_name']}: " f"{format_hover_number(pct_bandwidth, ',.2f')}% " f"({format_bandwidth(achieved_bandwidth)}/" f"{format_bandwidth(peak_bandwidth)})" diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 13247a79c226..9547efee3719 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -225,10 +225,33 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: assert unroofed[0]["points"][0]["hoverCells"] == [ "N/A", "N/A", - "Bandwidth:
 HBM: N/A% (900.000 GB/s/N/A)", + "Bandwidth:
\u2003HBM: N/A% (900.000 GB/s/N/A)", ] +def test_bandwidth_hover_shows_zero_levels() -> None: + """Every cache level still shows up at 0 when a kernel does no traffic + there, instead of vanishing from the tooltip.""" + ceiling = dict(CEILING, lds=[[0.01, 1.0], [1.0, 800.0], 800.0]) + _, model = kernel_traces( + make_roofline(["FP32"]), + { + "ai_hbm": [[1.0], [900.0]], + "ai_l2": [[0.0], [0.0]], + "ai_lds": [[0.0], [0.0]], + "kernelNames": ["kA"], + }, + sanitized_cache_hierarchy=["HBM", "L2", "LDS"], + ceiling_data=ceiling, + ) + bandwidth_html = model[0]["points"][0]["hoverCells"][2] + assert bandwidth_html == ( + "Bandwidth:
\u2003L2: N/A% (0.000 GB/s/N/A)" + "
\u2003HBM: 60.00% (900.000 GB/s/1.500 TB/s)" + "
\u2003LDS: 0.00% (0.000 GB/s/800.000 GB/s)" + ) + + def test_kernel_hover_carries_the_whole_name_up_to_the_limit() -> None: """A demangled name within the tooltip's length cap reaches the tooltip whole. It is wrapped onto as many lines as it takes, but nothing is From 4c641160dff6ddbcd6b93dc03f66937056bf1004 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 19:44:06 +0100 Subject: [PATCH 14/45] Proper performance limiting detection --- .../src/roofline/roofline_main.py | 22 ++++++++-------- .../tests/unit/roofline/test_roofline_main.py | 25 ++++++++++++++++--- 2 files changed, 33 insertions(+), 14 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index a35f066d671a..966d946c3a5c 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -372,14 +372,15 @@ def _determine_kernel_limiter( self, level_ai: dict[str, float], ceiling_data: dict[str, Any], - compute_cap: float, - compute_cap_label: str, + performance: float, + compute_peaks: list[tuple[str, float]], ) -> tuple[str, str]: """Name the specific binding roof for a kernel: the roof with the lowest - achievable performance at the kernel's operating point. The compute - candidate is the envelope cap the diagonals are actually drawn to, so - the limiter agrees with the drawn roof and with the percent of roofline - the tooltip reports. + achievable performance at the kernel's operating point. Each stacked + datatype's compute peak is judged on its own, and only peaks the + kernel's own measured performance doesn't already exceed are eligible + candidates, so the label names the specific datatype/op-path this + kernel is bound by instead of always the tallest peak drawn. Returns (label, category), where category is "Memory" or "Compute".""" candidates: list[tuple[float, str, str]] = [] @@ -388,8 +389,9 @@ def _determine_kernel_limiter( if bandwidth and ai_value > 0: candidates.append((bandwidth * ai_value, level_name, "Memory")) - if compute_cap != float("inf"): - candidates.append((compute_cap, compute_cap_label, "Compute")) + for label, peak in compute_peaks: + if peak >= performance: + candidates.append((peak, label, "Compute")) if not candidates: return "Unknown", "Unknown" @@ -414,7 +416,7 @@ def _build_kernel_traces( pct_runtime = self.__ai_data.get("pctRuntime", []) time_unit = self.__ai_data.get("timeUnit", "") total_dispatches = sum(count for count in counts if count is not None) - compute_cap, compute_cap_label = self._envelope_compute_cap(compute_peaks) + compute_cap, _ = self._envelope_compute_cap(compute_peaks) for kernel_index, kernel_name in enumerate(kernel_names): points, level_ai = self._build_kernel_points( @@ -431,7 +433,7 @@ def _build_kernel_traces( for values in (kernel_colors, counts, total_time, pct_runtime) ) limiter, limiter_category = self._determine_kernel_limiter( - level_ai, ceiling_data, compute_cap, compute_cap_label + level_ai, ceiling_data, points[0]["perf"], compute_peaks ) traces.append( diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 9547efee3719..30f868e7f9e4 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -182,9 +182,11 @@ def pct_roof(kernel: dict, point_index: int = 0) -> float: def test_kernel_traces_score_against_the_tallest_drawn_ceiling() -> None: - """A stacked figure caps points at the tallest compute roof drawn, so the - reported peak and limiter do not depend on the order datatypes were - stacked.""" + """A stacked figure caps every point's percent-of-roofline at the tallest + compute roof drawn, so that score does not depend on the order datatypes + were stacked -- but the limiter still only names a compute peak the + kernel's own performance does not exceed, falling back to the memory + level when no stacked peak is high enough to explain it.""" ai_data = {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]} matrix_traces, matrix_capped = kernel_traces(make_roofline(["FP32"]), ai_data) @@ -195,7 +197,22 @@ def test_kernel_traces_score_against_the_tallest_drawn_ceiling() -> None: assert pct_roof(matrix_capped[0]) < 100.0 assert pct_roof(valu_capped[0]) > 100.0 assert "Limited by Compute: FP32 MFMA" in matrix_traces[0].hovertemplate - assert "Limited by Compute: FP32 VALU" in valu_traces[0].hovertemplate + assert "Limited by Memory: HBM" in valu_traces[0].hovertemplate + + +def test_kernel_traces_name_the_specific_peak_not_the_tallest_stacked_one() -> None: + """When several compute peaks are stacked, the limiter names whichever one + actually bounds this kernel's own performance, not whichever peak happens + to be tallest on the figure.""" + ai_data = {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]} + + traces, _ = kernel_traces( + make_roofline(["FP32"]), + ai_data, + compute_peaks=[("FP32 MFMA", 90000.0), ("FP16 MFMA", 180000.0)], + ) + + assert "Limited by Compute: FP32 MFMA" in traces[0].hovertemplate def test_kernel_traces_name_the_roof_that_binds() -> None: From fdc6a6691faa37689a8558dfe4dbdb8158c1d0f9 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 20:55:52 +0100 Subject: [PATCH 15/45] Roofline performance value is calculated properly --- .../src/roofline/roofline_main.py | 83 +++++++------------ .../tests/unit/roofline/test_roofline_main.py | 63 ++++++++++++-- 2 files changed, 86 insertions(+), 60 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 966d946c3a5c..62a6aa3dd032 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -351,52 +351,39 @@ def _collect_compute_peaks(self, ops_flops: str) -> list[tuple[str, float]]: peaks.append((f"{dt} {label}", peak)) return peaks - def _roof_value_at( - self, - ai_value: float, - cache_key: str, - ceiling_data: dict[str, Any], - cap: float, - ) -> Optional[float]: - """Roofline throughput (peak) at this AI for the point's memory level: - min(bandwidth * AI, active compute cap); None when unavailable.""" - bandwidth = self._peak_value(ceiling_data, cache_key) - if not bandwidth or ai_value <= 0: - return None - roof = bandwidth * ai_value - if cap != float("inf"): - roof = min(roof, cap) - return roof if roof > 0 else None - def _determine_kernel_limiter( self, level_ai: dict[str, float], ceiling_data: dict[str, Any], performance: float, compute_peaks: list[tuple[str, float]], - ) -> tuple[str, str]: + ) -> tuple[str, str, Optional[float]]: """Name the specific binding roof for a kernel: the roof with the lowest achievable performance at the kernel's operating point. Each stacked - datatype's compute peak is judged on its own, and only peaks the - kernel's own measured performance doesn't already exceed are eligible - candidates, so the label names the specific datatype/op-path this - kernel is bound by instead of always the tallest peak drawn. - - Returns (label, category), where category is "Memory" or "Compute".""" + datatype's compute peak, and each memory level's own bandwidth * AI, + is judged on its own, and only candidates the kernel's own measured + performance doesn't already exceed are eligible, so the roof named (and + its value) is one this kernel could actually be bound by, instead of an + already-impossible ceiling or always the tallest peak drawn. + + Returns (label, category, roof_value); category is "Memory" or + "Compute", roof_value is None when no candidate is eligible.""" candidates: list[tuple[float, str, str]] = [] for level_name, ai_value in level_ai.items(): bandwidth = self._peak_value(ceiling_data, level_name.lower()) if bandwidth and ai_value > 0: - candidates.append((bandwidth * ai_value, level_name, "Memory")) + roof = bandwidth * ai_value + if roof >= performance: + candidates.append((roof, level_name, "Memory")) for label, peak in compute_peaks: if peak >= performance: candidates.append((peak, label, "Compute")) if not candidates: - return "Unknown", "Unknown" - _, label, category = min(candidates, key=lambda candidate: candidate[0]) - return label, category + return "Unknown", "Unknown", None + roof_value, label, category = min(candidates, key=lambda candidate: candidate[0]) + return label, category, roof_value def _build_kernel_traces( self, @@ -416,14 +403,12 @@ def _build_kernel_traces( pct_runtime = self.__ai_data.get("pctRuntime", []) time_unit = self.__ai_data.get("timeUnit", "") total_dispatches = sum(count for count in counts if count is not None) - compute_cap, _ = self._envelope_compute_cap(compute_peaks) for kernel_index, kernel_name in enumerate(kernel_names): - points, level_ai = self._build_kernel_points( + points, level_ai, bandwidth_html = self._build_kernel_points( kernel_index=kernel_index, sanitized_cache_hierarchy=sanitized_cache_hierarchy, ceiling_data=ceiling_data, - compute_cap=compute_cap, ) if not points: continue @@ -432,9 +417,19 @@ def _build_kernel_traces( values[kernel_index] if kernel_index < len(values) else None for values in (kernel_colors, counts, total_time, pct_runtime) ) - limiter, limiter_category = self._determine_kernel_limiter( + limiter, limiter_category, roof_value = self._determine_kernel_limiter( level_ai, ceiling_data, points[0]["perf"], compute_peaks ) + pct_roof = ( + 100.0 * points[0]["perf"] / roof_value if roof_value else None + ) + hover_cells = [ + format_hover_number(roof_value, ",.0f"), + format_hover_number(pct_roof, ",.2f"), + bandwidth_html, + ] + for point in points: + point["hoverCells"] = hover_cells traces.append( go.Scatter( @@ -476,11 +471,11 @@ def _build_kernel_points( kernel_index: int, sanitized_cache_hierarchy: list[str], ceiling_data: dict[str, Any], - compute_cap: float, - ) -> tuple[list[dict[str, Any]], dict[str, float]]: + ) -> tuple[list[dict[str, Any]], dict[str, float], str]: """One kernel's plotted points, one per memory level it has data for. - Also returns the level -> AI map the limiter is chosen from. + Also returns the level -> AI map the limiter is chosen from, and the + 'Bandwidth:' hover block shared by every one of this kernel's points. """ points: list[dict[str, Any]] = [] bandwidth_entries: list[dict[str, Any]] = [] @@ -510,22 +505,13 @@ def _build_kernel_points( }) continue - roof_perf = self._roof_value_at( - ai_value=ai_value, - cache_key=cache_key, - ceiling_data=ceiling_data, - cap=compute_cap, - ) # This kernel's own achieved bandwidth at this level (not the # hardware ceiling): performance (FLOP/s) / AI (FLOP/Byte) = Byte/s. achieved_bandwidth = performance / ai_value - pct_roof = 100.0 * performance / roof_perf if roof_perf else None points.append({ "peak": level_name, "ai": ai_value, "perf": performance, - "roof_perf_txt": format_hover_number(roof_perf, ",.0f"), - "pct_roof_txt": format_hover_number(pct_roof, ",.2f"), }) bandwidth_entries.append({ "level_name": level_name, @@ -535,14 +521,7 @@ def _build_kernel_points( level_ai[level_name] = ai_value bandwidth_html = self._build_bandwidth_hover_html(bandwidth_entries) - for point in points: - point["hoverCells"] = [ - point.pop("roof_perf_txt"), - point.pop("pct_roof_txt"), - bandwidth_html, - ] - - return points, level_ai + return points, level_ai, bandwidth_html @staticmethod def _build_bandwidth_hover_html(bandwidth_entries: list[dict[str, Any]]) -> str: diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 30f868e7f9e4..59944f931133 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -181,12 +181,14 @@ def pct_roof(kernel: dict, point_index: int = 0) -> float: return float(kernel["points"][point_index]["hoverCells"][1]) -def test_kernel_traces_score_against_the_tallest_drawn_ceiling() -> None: - """A stacked figure caps every point's percent-of-roofline at the tallest - compute roof drawn, so that score does not depend on the order datatypes - were stacked -- but the limiter still only names a compute peak the - kernel's own performance does not exceed, falling back to the memory - level when no stacked peak is high enough to explain it.""" +def test_kernel_traces_score_against_the_roof_that_binds_not_the_tallest_drawn() -> None: + """The percentage is scored against whichever roof the limiter actually + names, not always the tallest compute roof stacked on the figure: dropping + the tall MFMA peak from the candidate set changes both the label (falls + back to the memory level, since VALU's peak is below this kernel's own + performance) and the percentage (scored against HBM's own 150,000 + GFLOP/s implied ceiling instead of MFMA's 90,000), and neither case's + percentage exceeds 100%.""" ai_data = {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]} matrix_traces, matrix_capped = kernel_traces(make_roofline(["FP32"]), ai_data) @@ -194,8 +196,8 @@ def test_kernel_traces_score_against_the_tallest_drawn_ceiling() -> None: make_roofline(["FP32"]), ai_data, compute_peaks=[("FP32 VALU", 9000.0)] ) - assert pct_roof(matrix_capped[0]) < 100.0 - assert pct_roof(valu_capped[0]) > 100.0 + assert pct_roof(matrix_capped[0]) == pytest.approx(55.56, abs=0.01) + assert pct_roof(valu_capped[0]) == pytest.approx(33.33, abs=0.01) assert "Limited by Compute: FP32 MFMA" in matrix_traces[0].hovertemplate assert "Limited by Memory: HBM" in valu_traces[0].hovertemplate @@ -215,6 +217,22 @@ def test_kernel_traces_name_the_specific_peak_not_the_tallest_stacked_one() -> N assert "Limited by Compute: FP32 MFMA" in traces[0].hovertemplate +def test_kernel_traces_ignore_a_memory_roof_the_kernel_already_exceeds() -> None: + """A memory level's implied ceiling (bandwidth * AI) that the kernel's own + measured performance already exceeds is not an achievable roof, so the + limiter must not name it -- exactly as compute peaks below performance are + already excluded. Here HBM implies a 1,500 GFLOP/s ceiling but the kernel + measures 5,000 GFLOP/s, so the limiter must fall through to the cheapest + eligible compute peak instead of naming the impossible memory ceiling.""" + traces, _ = kernel_traces( + make_roofline(["FP32"]), + {"ai_hbm": [[1.0], [5000.0]], "kernelNames": ["kA"]}, + ) + + assert "Limited by Compute: FP32 VALU" in traces[0].hovertemplate + assert "Limited by Memory: HBM" not in traces[0].hovertemplate + + def test_kernel_traces_name_the_roof_that_binds() -> None: """A kernel whose bandwidth roof sits under the compute cap is limited by its memory level, and falls back to Unknown when the ceiling data holds no roof @@ -246,6 +264,35 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: ] +def test_kernel_traces_share_one_performance_value_across_a_kernels_points() -> None: + """A kernel's HBM point and L2 point measure the same achieved performance + but sit at different arithmetic intensities, so today each point is scored + against its own level's bandwidth -- 25.00% at HBM, 66.67% at L2. Once the + limiter's own closest-roof search (L2's 1,500 GFLOP/s implied ceiling beats + HBM's 4,000) drives the displayed percentage too, both of the kernel's + points must show that same 66.67% / 1,500, not their own per-level value.""" + ceiling = { + "hbm": [[0.01, 1.0], [1.0, 2000.0], 2000.0], + "l2": [[0.01, 1.0], [1.0, 3000.0], 3000.0], + } + + _, model = kernel_traces( + make_roofline(["FP32"]), + { + "ai_hbm": [[2.0], [1000.0]], + "ai_l2": [[0.5], [1000.0]], + "kernelNames": ["kA"], + }, + sanitized_cache_hierarchy=["HBM", "L2"], + ceiling_data=ceiling, + compute_peaks=[], + ) + + points = {point["peak"]: point for point in model[0]["points"]} + assert points["HBM"]["hoverCells"][:2] == ["1,500", "66.67"] + assert points["L2"]["hoverCells"][:2] == ["1,500", "66.67"] + + def test_bandwidth_hover_shows_zero_levels() -> None: """Every cache level still shows up at 0 when a kernel does no traffic there, instead of vanishing from the tooltip.""" From 230976c3681aab4095bbfb5a8e8a810f53d297e7 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 21:00:41 +0100 Subject: [PATCH 16/45] Fix for bandwidth parameters --- .../src/roofline/roofline_hover.py | 64 +++++++++++++------ .../src/roofline/roofline_main.py | 5 +- 2 files changed, 45 insertions(+), 24 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index 4d44f849d01f..30a415cc295a 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -69,7 +69,7 @@ def build_kernel_hover_template( f"Limited by {limiter_category}: {limiter}", f"Performance: %{{customdata[1]}}% (%{{y:,.0f}} / %{{customdata[0]}} {unit})", "%{customdata[2]}", - f"Dispatch Number: {pct_dispatches_txt}% ({_format_integer(count)} / {_format_integer(total_dispatches)})", + f"Dispatch Count: {pct_dispatches_txt}% ({_format_integer(count)} / {_format_integer(total_dispatches)})", f"Duration: {pct_runtime_txt}% ({kernel_time_txt} / {total_app_time_txt} {time_unit})", "AI: %{x:.6g}", ], @@ -110,24 +110,26 @@ def build_compute_peak_hover( ) -def format_hover_number(value: object, spec: str) -> str: - """Format a numeric tooltip value with the given format spec, or N/A.""" +def _safe_float(value: object) -> Optional[float]: + """Convert to float, or None if missing/not numeric.""" if value is None: - return "N/A" + return None try: - return format(float(value), spec) + return float(value) except (TypeError, ValueError): - return "N/A" + return None + + +def format_hover_number(value: object, spec: str) -> str: + """Format a numeric tooltip value with the given format spec, or N/A.""" + parsed = _safe_float(value) + return format(parsed, spec) if parsed is not None else "N/A" def _format_integer(value: object) -> str: """Thousands-separated integer for the tooltip, or N/A when missing.""" - if value is None: - return "N/A" - try: - return f"{int(round(float(value))):,}" - except (TypeError, ValueError): - return "N/A" + parsed = _safe_float(value) + return f"{int(round(parsed)):,}" if parsed is not None else "N/A" def _hover(header: str, rows: list[str]) -> str: @@ -135,13 +137,33 @@ def _hover(header: str, rows: list[str]) -> str: return "
".join([header, "", *rows]) + "" -def format_bandwidth(gb_per_s: float) -> str: - """Bandwidth as GB/s, switching to TB/s at >= 1000 GB/s - so the roof hover stays readable.""" - try: - value = float(gb_per_s) - except (TypeError, ValueError): - return "N/A" +def _bandwidth_unit(value: float) -> tuple[float, str]: + """Scale factor and suffix for a bandwidth magnitude: GB/s, switching to + TB/s at >= 1000 GB/s so the hover stays readable.""" if abs(value) >= 1000.0: - return f"{value / 1000.0:,.3f} TB/s" - return f"{value:,.3f} GB/s" + return 1000.0, "TB/s" + return 1.0, "GB/s" + + +def format_bandwidth(gb_per_s: object) -> str: + """Bandwidth as GB/s, switching to TB/s at >= 1000 GB/s.""" + value = _safe_float(gb_per_s) + if value is None: + return "N/A" + scale, unit = _bandwidth_unit(value) + return f"{value / scale:,.3f} {unit}" + + +def format_bandwidth_pair(achieved_gb_per_s: object, peak_gb_per_s: object) -> str: + """Achieved/peak bandwidth sharing a single unit suffix (GB/s or TB/s), + picked from the peak's magnitude, matching the single-suffix style of + the Performance field instead of scaling each side independently.""" + achieved = _safe_float(achieved_gb_per_s) + peak = _safe_float(peak_gb_per_s) + reference = peak if peak else achieved + scale, unit = _bandwidth_unit(reference) if reference is not None else (1.0, "GB/s") + achieved_txt = format_hover_number( + achieved / scale if achieved is not None else None, ",.3f" + ) + peak_txt = format_hover_number(peak / scale if peak is not None else None, ",.3f") + return f"{achieved_txt} / {peak_txt} {unit}" diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 62a6aa3dd032..e54c6a56dda6 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -17,7 +17,7 @@ build_compute_peak_hover, build_kernel_hover_template, build_roof_hover, - format_bandwidth, + format_bandwidth_pair, format_hover_number, truncate_kernel_name, wrap_hover_name, @@ -543,8 +543,7 @@ def _build_bandwidth_hover_html(bandwidth_entries: list[dict[str, Any]]) -> str: lines.append( f"\u2003{entry['level_name']}: " f"{format_hover_number(pct_bandwidth, ',.2f')}% " - f"({format_bandwidth(achieved_bandwidth)}/" - f"{format_bandwidth(peak_bandwidth)})" + f"({format_bandwidth_pair(achieved_bandwidth, peak_bandwidth)})" ) return "
".join(lines) From d407a28e6155038338fc793a7ea4f40d8a4540fb Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 21:05:42 +0100 Subject: [PATCH 17/45] hover formatting --- projects/rocprofiler-compute/src/roofline/roofline_hover.py | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index 30a415cc295a..b63f649c1520 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -68,10 +68,12 @@ def build_kernel_hover_template( [ f"Limited by {limiter_category}: {limiter}", f"Performance: %{{customdata[1]}}% (%{{y:,.0f}} / %{{customdata[0]}} {unit})", + "AI: %{x:.6g}", + "", + "Details", "%{customdata[2]}", f"Dispatch Count: {pct_dispatches_txt}% ({_format_integer(count)} / {_format_integer(total_dispatches)})", f"Duration: {pct_runtime_txt}% ({kernel_time_txt} / {total_app_time_txt} {time_unit})", - "AI: %{x:.6g}", ], ) From 3b2b0c376350d27098a8ec76c69c2752a3b5fd55 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Wed, 23 Sep 2026 21:43:23 +0100 Subject: [PATCH 18/45] Option to visualize kernels by limiting peak --- .../specs/roofline-limiting-ai-axis-option.md | 72 +++++++++++++++++++ .../src/roofline/assets/roofline_plot.js | 21 +++++- .../src/roofline/roofline_html.py | 5 +- .../src/roofline/roofline_main.py | 6 ++ .../tests/unit/roofline/test_roofline_main.py | 64 +++++++++++++++-- 5 files changed, 160 insertions(+), 8 deletions(-) create mode 100644 projects/rocprofiler-compute/specs/roofline-limiting-ai-axis-option.md diff --git a/projects/rocprofiler-compute/specs/roofline-limiting-ai-axis-option.md b/projects/rocprofiler-compute/specs/roofline-limiting-ai-axis-option.md new file mode 100644 index 000000000000..65cecdcbd008 --- /dev/null +++ b/projects/rocprofiler-compute/specs/roofline-limiting-ai-axis-option.md @@ -0,0 +1,72 @@ +# Roofline: "Limiting" option for the AI axis dropdown + +## Problem Statement + +In the interactive roofline HTML report, each kernel can appear as multiple dots on the chart — one per memory (cache) level it has traffic for (e.g. L0, L1, L2, HBM, LDS), each plotted at that level's own arithmetic intensity (AI) with the kernel's performance. The "AI axis" dropdown lets a user pick one specific level to display, or "All peaks" to show every level's dot for every kernel at once. + +There is no option that shows, for every kernel simultaneously, only the one dot that actually reflects its bottleneck. A user who wants to scan the whole chart for "what's actually limiting each kernel" today has to either pick a single memory level (which may not be the level that kernel is actually bound by) or view "All peaks" and mentally filter the clutter of every kernel's non-limiting dots. This mirrors the trade-off between Intel Advisor's single-dot CARM view and its multi-dot Integrated/Memory-Level roofline view, discussed as prior art earlier in this workstream. + +## Solution + +Add a new `"Limiting"` entry to the existing "AI axis" dropdown. When selected, every kernel is drawn with exactly one dot: the point corresponding to whichever memory level (or, for compute-bound/unknown kernels, a well-defined fallback point) the backend has already identified as that kernel's limiter. This reuses the limiter/limiter_category detection already computed for the hover tooltip ("Limited by ...") — it is not new bottleneck-detection logic, only a new way to visualize the existing determination. + +## User Stories + +1. As a performance engineer scanning a roofline chart with many kernels, I want an option to show only each kernel's limiting dot, so that I can see at a glance which kernels are bandwidth-bound vs. compute-bound without visual clutter from non-limiting cache levels. +2. As a performance engineer, I want the "Limiting" option to sit in the same "AI axis" dropdown I already use to switch between cache levels, so that I don't have to learn a new control. +3. As a performance engineer, I want the existing per-level options and "All peaks" to keep working exactly as before, so that adding this option doesn't change behavior I already rely on. +4. As a performance engineer, I want the default dropdown selection to remain unchanged (still the current default level, e.g. HBM if present), so that opening a report I've already generated doesn't silently change its appearance. +5. As a performance engineer, I want a memory-bound kernel's lone dot in "Limiting" mode to sit exactly at that cache level's own AI/performance point, so that its horizontal position still means what it means everywhere else on the chart. +6. As a performance engineer, I want a compute-bound kernel (no single cache level is its bottleneck) to still show exactly one dot rather than vanishing from the chart, so that I don't lose visibility into compute-bound kernels while in this mode. +7. As a performance engineer, I want that fallback dot (for compute-bound or "Unknown"-limiter kernels) to be the leftmost (lowest-AI) of the kernel's available points, so that the fallback is deterministic and documented rather than arbitrary. +8. As a performance engineer, I want the dot's color to still reflect its cache level (the same color coding used everywhere else, including in single-kernel isolation view), so that color continues to mean the same thing across every view of the chart. +9. As a performance engineer, I want the hover tooltip on a "Limiting"-mode dot to be unchanged (still "Limited by {category}: {name}", Performance, AI, Details/Bandwidth/Dispatch/Duration), so that the information I already rely on in hover text isn't disrupted by this new mode. +10. As a performance engineer, I want to click/isolate a single kernel while in "Limiting" mode and see all of that kernel's cache-level dots (not just the limiting one), so that isolating a kernel still gives me the full multi-level picture, consistent with how isolation already overrides the other specific-level dropdown options. +11. As a performance engineer, I want the "Limiting" mode to leave the bandwidth-roof lines and compute-peak (flat roof) lines completely untouched, so that the roofline geometry I compare dots against never changes based on the AI-axis selection. +12. As a performance engineer, I want "Limiting" mode to apply consistently across the whole report regardless of which datatypes are currently toggled visible via the separate "Precision" control, so that the two controls compose predictably instead of interacting in surprising ways. +13. As a developer maintaining this codebase, I want the "which single point is limiting" decision made once, server-side, in Python, so that the client-side JavaScript only needs a trivial lookup rather than re-implementing bottleneck-detection logic. +14. As a developer maintaining this codebase, I want the new field driving this decision to be unit-testable via the existing kernel-trace-building tests, so that regressions in limiter-to-dot mapping are caught the same way regressions in the existing limiter/hover logic already are. + +## Implementation Decisions + +- **New dropdown option**: add `"Limiting"` as an entry in the "AI axis" `` (`roofline-peak-select`), positioned first in the option list (before the per-level entries, before "All peaks"). It is additive: existing per-level options, the `"All peaks"` option, and the current default selection logic are unchanged. -- **New special peak value**: introduce a new sentinel value (analogous to the existing `ALL_PEAKS_VALUE = "all"` sentinel) to represent this mode, e.g. a `LIMITING_PEAK_VALUE` constant, threaded through the same places `ALL_PEAKS_VALUE` currently is (model construction in Python, `buildPeakOptions()`/`effectivePeak()`/`pointsForCurrentPeak()` in JS). -- **Server-side resolution of the limiting point** (the core of this feature): in `_build_kernel_traces`, at the point where `_determine_kernel_limiter` is already called once per kernel (using `level_ai` and `compute_peaks`), resolve its result to a concrete point selector and add it to the kernel's model as a new plain field (e.g. `limitingPeak`), holding a value that matches one of that kernel's own `point["peak"]` values: - - If `limiter_category == "Memory"`: `limitingPeak` = the limiter's cache-level name (already directly equal to one of the kernel's `point["peak"]` values, by construction of `_determine_kernel_limiter`). - - If `limiter_category` is `"Compute"` or `"Unknown"`: `limitingPeak` = the `peak` of the kernel's leftmost point (the point with the minimum `ai` value among its own `points` list). -- **No change to limiter detection itself**: `_determine_kernel_limiter`'s algorithm (eligible candidates, minimum-roof selection, `("Unknown", "Unknown", None)` fallback) is unchanged. This feature only adds a downstream mapping from its existing result to a displayable point. -- **Client-side filtering**: `pointsForCurrentPeak(kernel)` (or equivalent) gains one more branch: when the effective peak is the new sentinel, return only the point whose `peak === kernel.limitingPeak`, instead of all points or a single specific-level match. No bottleneck logic is duplicated client-side — it's a direct equality filter against the pre-resolved field. -- **Kernel isolation unaffected**: `effectivePeak()`'s existing override (forcing `"all"` whenever exactly one kernel is isolated, regardless of the dropdown's value) is not special-cased for the new sentinel — it already applies uniformly to every non-`"all"` dropdown value, including this new one. -- **Coloring unaffected**: dot coloring continues to key off `point.peak` (`peakColors[point.peak]`) exactly as it does today for every other specific-level selection and for isolation view; no new color/style branch is introduced for the compute/unknown fallback case. -- **No change to roofline line rendering**: bandwidth-roof and compute-peak line traces are entirely independent of the AI-axis selection today and remain so. -- **No change to hover template**: `build_kernel_hover_template` and its inputs (`limiter`, `limiter_category`, per-kernel shared hover cells) are unaffected; the existing "Limited by ..." text already reflects the same underlying determination this feature now also uses to pick a dot. -- **Single combined figure**: confirmed the report renders one combined Plotly figure with one shared "AI axis" dropdown/state (Ops vs. Flops are trace subsets toggled by the separate "Precision" control, not two figures). No per-figure scoping is needed; the new option applies uniformly to whatever kernel traces are currently present. -- **Tooltip copy**: the informational tooltip text next to the "AI axis" label (`PEAK_TITLE`) should be updated to mention the new option, at the implementer's discretion for exact wording — not a decision point requiring further sign-off. - -## Testing Decisions - -- Tests should assert on externally observable output — the kernel `model` dict returned by the kernel-trace-building code path (`points`, `limitingPeak`, `hoverCells`) and/or the rendered `hovertemplate` string — not on internal call sequencing. -- **Primary seam**: the existing `kernel_traces()` test helper in `tests/unit/roofline/test_roofline_main.py`, which already wraps `_build_kernel_traces` and is used by the existing limiter tests (`test_kernel_traces_score_against_the_roof_that_binds_not_the_tallest_drawn`, `test_kernel_traces_name_the_specific_peak_not_the_tallest_stacked_one`, `test_kernel_traces_ignore_a_memory_roof_the_kernel_already_exceeds`, `test_kernel_traces_name_the_roof_that_binds`, `test_kernel_traces_share_one_performance_value_across_a_kernels_points`). New tests should follow the same pattern and assert the new `limitingPeak` field, covering: - - A memory-bound kernel: `limitingPeak` equals the cache level `_determine_kernel_limiter` names. - - A compute-bound kernel: `limitingPeak` equals the `peak` of the point with the lowest `ai` value. - - An "Unknown"-limiter kernel (no ceiling/compute data, mirroring the existing "Unknown"/"Unknown" test case): `limitingPeak` falls back the same way as the compute-bound case. - - A kernel with a single available point: `limitingPeak` trivially equals that point's `peak` regardless of category. -- **No JS test infrastructure exists in this repo** (no jest/vitest/node test runner; `roofline_plot.js` is currently only indirectly covered by Python tests that regex-match against the JS source text, e.g. `test_the_controller_looks_up_controls_the_page_renders`). This spec deliberately keeps the decision logic server-side so the client-side change (`point.peak === kernel.limitingPeak`) is trivial enough not to require introducing new JS test infrastructure. If a maintainer judges dedicated JS coverage is warranted anyway, that's a follow-up decision outside this spec, not a blocker for it. -- Existing tests asserting the current dropdown options, default peak selection, and "All peaks" behavior must continue to pass unmodified — this feature must not alter any of that existing behavior. - -## Out of Scope - -- Any change to how `_determine_kernel_limiter` decides what counts as a candidate roof/peak, or how it picks among candidates. -- Any change to bandwidth-roof or compute-peak line rendering. -- Any change to the "Precision" dropdown or its interaction model with dataset filtering. -- Any change to the hover tooltip content or ordering (already settled in prior work: Limited by → Performance → AI → Details → Bandwidth/Dispatch Count/Duration). -- A legend or on-chart visual distinction between "this dot is the true limiter" vs. "this dot is a leftmost fallback" for compute-bound/unknown kernels — the fallback is silent (no special styling), per the confirmed design. -- Introducing JS test infrastructure (jest/vitest/node) to this repo. -- Any change to per-kernel isolation's existing behavior of always showing all cache-level dots. -- Publishing this spec to an external issue tracker — per explicit request, this spec is written to a local markdown file instead. - -## Further Notes - -- This feature is intentionally a pure "visualization of an existing determination," not new bottleneck-detection logic — keeping the surface area small was an explicit design goal reached during the interview that produced this spec. -- The design was validated against Intel Advisor's CARM (single aggregate dot) vs. Integrated/Memory-Level Roofline (one dot per cache level) distinction as prior art; "Limiting" mode is closer in spirit to Advisor's Integrated model filtered down to just the binding level, rather than to CARM's L1-anchored cumulative-traffic AI (this codebase's per-level AI values are not being replaced or recomputed — only filtered down to one per kernel). -- The exact constant name (`LIMITING_PEAK_VALUE` or similar) and the new model field name (`limitingPeak` or similar) are naming suggestions, not binding decisions — an implementer should follow this codebase's existing naming conventions when landing them. From 2f6b27415a81ac931dd8c1874d139daa1f12540e Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 14:50:07 +0100 Subject: [PATCH 21/45] removed dev script --- .../run_rocprof_analysis.sh | 30 ------------------- 1 file changed, 30 deletions(-) delete mode 100755 projects/rocprofiler-compute/run_rocprof_analysis.sh diff --git a/projects/rocprofiler-compute/run_rocprof_analysis.sh b/projects/rocprofiler-compute/run_rocprof_analysis.sh deleted file mode 100755 index b87e9bb3fb68..000000000000 --- a/projects/rocprofiler-compute/run_rocprof_analysis.sh +++ /dev/null @@ -1,30 +0,0 @@ -#!/bin/bash - -set -euo pipefail - -if [ $# -ne 1 ]; then - echo "Usage: $0 " - exit 1 -fi - -FOLDER="$1" - -if [ ! -d "$FOLDER" ]; then - echo "Error: '$FOLDER' is not a directory" - exit 1 -fi - -for full_path in "$FOLDER"/*; do - [ -f "$full_path" ] || continue - - file_name="$(basename "$full_path")" - - echo "=== Profiling: $file_name ===" - ./src/rocprof-compute profile --roof-only \ - --output-directory "/work-space/output/bmm/$file_name" \ - --overwrite \ - -- python "$full_path" - - echo "=== Analyzing: $file_name ===" - ./src/rocprof-compute analyze -p "/work-space/output/bmm/$file_name" -done From 83a96b6721b55528e2eb3a20dfa633d24506250a Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 14:55:09 +0100 Subject: [PATCH 22/45] roofline peak instead of average --- .../src/roofline/benchmark/benchmark_base.py | 165 ++++++------------ 1 file changed, 56 insertions(+), 109 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py index 08bc3b3d2743..4e1efb05ca16 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py @@ -9,9 +9,7 @@ # ----------------------------------------------------------------------------- import csv -import math from abc import ABC -from collections import namedtuple from collections.abc import Generator from contextlib import contextmanager from ctypes import ( @@ -39,9 +37,6 @@ # GLOBAL VARIABLES # ============================================================================= -Stats = namedtuple("Stats", ["mean", "stdev", "confidence"]) -PerfMetrics = namedtuple("PerfMetrics", ["mean", "low", "high"]) - DEFAULT_WORKGROUP_SIZE = 256 DEFAULT_WORKGROUPS = 8192 DEFAULT_THREADS = DEFAULT_WORKGROUP_SIZE * DEFAULT_WORKGROUPS @@ -71,8 +66,6 @@ def __init__(self, device_id: int, cache_sizes: dict) -> None: self.csv_cols_map: dict[str, str] self.WAVEFRONT_SIZE: int self.MATRIX_OPS_TYPE: str - self.event_start = hip.hipEventCreate() - self.event_stop = hip.hipEventCreate() self.cache_sizes = cache_sizes @@ -169,19 +162,9 @@ def show_progress(self, pct: float) -> None: print(f"\r{int(pct * 100):3d}% {bar}", end="", flush=True) - def calc_stats(self, samples: list) -> Stats: - """Returns a named tuple with the mean, std deviation and confidence.""" - mean = sum(samples) / len(samples) - m = max(samples) - - stdev = 0.0 - - for i in range(len(samples)): - stdev += math.pow(samples[i] - mean, 2) - - stdev = math.sqrt(stdev / len(samples)) - - return Stats(m, stdev, 1.96 * stdev / math.sqrt(len(samples))) + def calc_stats(self, samples: list) -> float: + """Returns the peak (max) value from the samples.""" + return max(samples) class Program: """Helper class for loading and compiling kernels.""" @@ -416,7 +399,7 @@ def set_kernel_source(self) -> None: # ----------------------------------------------------------------------------- # HBM bandwidth benchmark - def hbm_bw_benchmark(self, device: int) -> PerfMetrics: + def hbm_bw_benchmark(self, device: int) -> float: num_experiments = DEFAULT_NUM_EXPERIMENTS hip.hipSetDevice(device) @@ -463,28 +446,21 @@ def hbm_bw_benchmark(self, device: int) -> PerfMetrics: [d_src, c_int64(num_steps)], ) - stats = self.calc_stats(samples) - - mean = stats.mean - stdev = stats.stdev - - perf_metrics = PerfMetrics( - mean, mean - stats.confidence, mean + stats.confidence - ) + peak = self.calc_stats(samples) - event_ms = total_bytes / mean / 1e6 + event_ms = total_bytes / peak / 1e6 print( f"HBM BW, GPU ID: {device}, workgroupSize:{workgroup_size}, " f"workgroups:{workgroups}, experiments:{num_experiments}, " f"traffic:{total_bytes} bytes, duration:{event_ms:.1f} ms, " - f"mean:{mean:.1f} GB/sec, stdev:{stdev:.1f} GB/sec" + f"peak:{peak:.1f} GB/sec" ) - return perf_metrics + return peak # Generic cache bandwidth benchmark - def cache_bw_bench(self, device: int, type: str, iters: int) -> PerfMetrics: + def cache_bw_bench(self, device: int, type: str, iters: int) -> float: hip.hipSetDevice(device) num_experiments = DEFAULT_NUM_EXPERIMENTS @@ -525,43 +501,37 @@ def cache_bw_bench(self, device: int, type: str, iters: int) -> PerfMetrics: [mem_block, dummy, iters], ) - stats = self.calc_stats(samples) - mean = stats.mean - stdev = stats.stdev + peak = self.calc_stats(samples) - perf_metrics = PerfMetrics( - mean, mean - stats.confidence, mean + stats.confidence - ) - - event_ms = total_bytes / mean / 1e6 + event_ms = total_bytes / peak / 1e6 print( f"{type} BW, GPU ID: {device}, workgroupSize:{workgroup_size}, " f"workgroups:{workgroups}, experiments:{num_experiments}, " f"traffic:{total_bytes} bytes, duration:{event_ms:.1f} ms, " - f"mean:{mean:.1f} GB/sec, stdev:{stdev:1f} GB/sec" + f"peak:{peak:.1f} GB/sec" ) - return perf_metrics + return peak # MALL cache bandwidth benchmark - def mall_bw_bench(self, device: int) -> PerfMetrics: + def mall_bw_bench(self, device: int) -> float: return self.cache_bw_bench(device, "MALL", 1) # L0 cache bandwidth benchmark - def l0_bw_bench(self, device: int) -> PerfMetrics: + def l0_bw_bench(self, device: int) -> float: return self.cache_bw_bench(device, "L0", 100) # L1 cache bandwidth benchmark - def l1_bw_bench(self, device: int) -> PerfMetrics: + def l1_bw_bench(self, device: int) -> float: return self.cache_bw_bench(device, "L1", 100) # L2 cache bandwidth benchmark - def l2_bw_bench(self, device: int) -> PerfMetrics: + def l2_bw_bench(self, device: int) -> float: return self.cache_bw_bench(device, "L2", 10) # LDS cache bandwidth benchmark - def lds_bw_benchmark(self, device: int) -> PerfMetrics: + def lds_bw_benchmark(self, device: int) -> float: num_experiments = DEFAULT_NUM_EXPERIMENTS workgroup_size = DEFAULT_WORKGROUP_SIZE @@ -594,27 +564,21 @@ def lds_bw_benchmark(self, device: int) -> PerfMetrics: [iters, dummy], ) - stats = self.calc_stats(samples) - mean = stats.mean - stdev = stats.stdev - - perf_metrics = PerfMetrics( - mean, mean - stats.confidence, mean + stats.confidence - ) + peak = self.calc_stats(samples) - event_ms = total_bytes / mean / 1e6 + event_ms = total_bytes / peak / 1e6 print( f"LDS BW, GPU ID: {device}, workgroupSize:{workgroup_size}, " f"workgroups:{workgroups}, experiments:{num_experiments}, " f"traffic:{total_bytes} bytes, duration:{event_ms:.1f} ms, " - f"mean:{mean:.1f} GB/sec, stdev:{stdev:1f} GB/sec" + f"peak:{peak:.1f} GB/sec" ) - return perf_metrics + return peak # Generic FLOPs benchmark - def flops_bench(self, device: int, type: str, unit: str, rate: int) -> PerfMetrics: + def flops_bench(self, device: int, type: str, unit: str, rate: int) -> float: num_experiments = DEFAULT_NUM_EXPERIMENTS workgroup_size = DEFAULT_WORKGROUP_SIZE cus = hip.hipGetDeviceProperties(device).multiProcessorCount @@ -663,28 +627,22 @@ def flops_bench(self, device: int, type: str, unit: str, rate: int) -> PerfMetri [memblock, iterations], ) - stats = self.calc_stats(samples) - mean = stats.mean - stdev = stats.stdev + peak = self.calc_stats(samples) - perf_metrics = PerfMetrics( - mean, mean - stats.confidence, mean + stats.confidence - ) - - event_ms = total_flops / mean / 1e6 + event_ms = total_flops / peak / 1e6 print( f"Peak VALU {unit}s ({type}), GPU ID: {device}, " f"workgroupSize:{workgroup_size}, " f"workgroups:{workgroups}, experiments:{num_experiments}, " f"{unit}:{total_flops}, duration:{event_ms:.1f} ms, " - f"mean:{mean:.1f} {rate}, stdev={stdev:.1f} GFLOPS" + f"peak:{peak:.1f} {rate}" ) - return perf_metrics + return peak # Generic matrix operations benchmark - def matrix_bench(self, device: int, type: str, unit: str, rate: int) -> PerfMetrics: + def matrix_bench(self, device: int, type: str, unit: str, rate: int) -> float: experiments = DEFAULT_NUM_EXPERIMENTS iters = 2000 @@ -734,81 +692,74 @@ def matrix_bench(self, device: int, type: str, unit: str, rate: int) -> PerfMetr [iters, dummy], ) - stats = self.calc_stats(samples) - mean = stats.mean - stdev = stats.stdev - - perf_metrics = PerfMetrics( - mean, mean - stats.confidence, mean + stats.confidence - ) + peak = self.calc_stats(samples) - event_ms = total_flops / mean / 1e6 + event_ms = total_flops / peak / 1e6 print( f"Peak {self.MATRIX_OPS_TYPE} {unit}s ({type}), GPU ID: {device}, " f"workgroupSize:{workgroup_size}, workgroups:{workgroups}, " f"experiments:{experiments}, {unit}:{total_flops}, " - f"duration:{event_ms:.2f} ms, mean:{mean:.1f} {rate}, " - f"stdev:{stdev:.1f} GFLOPS" + f"duration:{event_ms:.2f} ms, peak:{peak:.1f} {rate}" ) - return perf_metrics + return peak - def matrix_f32_bench(self, device: int) -> PerfMetrics: + def matrix_f32_bench(self, device: int) -> float: return self.matrix_bench(device, "F32", "FLOP", "GFLOPS") - def matrix_f16_bench(self, device: int) -> PerfMetrics: + def matrix_f16_bench(self, device: int) -> float: return self.matrix_bench(device, "F16", "FLOP", "GFLOPS") - def matrix_bf16_bench(self, device: int) -> PerfMetrics: + def matrix_bf16_bench(self, device: int) -> float: return self.matrix_bench(device, "BF16", "FLOP", "GFLOPS") - def matrix_f64_bench(self, device: int) -> PerfMetrics: + def matrix_f64_bench(self, device: int) -> float: return self.matrix_bench(device, "F64", "FLOP", "GFLOPS") - def matrix_f8_bench(self, device: int) -> PerfMetrics: + def matrix_f8_bench(self, device: int) -> float: return self.matrix_bench(device, "F8", "FLOP", "GFLOPS") - def matrix_i8_bench(self, device: int) -> PerfMetrics: + def matrix_i8_bench(self, device: int) -> float: return self.matrix_bench(device, "I8", "IOP", "GOPS") - def matrix_f4_bench(self, device: int) -> PerfMetrics: + def matrix_f4_bench(self, device: int) -> float: return self.matrix_bench(device, "F4", "FLOP", "GFLOPS") - def matrix_f6_bench(self, device: int) -> PerfMetrics: + def matrix_f6_bench(self, device: int) -> float: return self.matrix_bench(device, "F6", "FLOP", "GFLOPS") - def matrix_f6f4_bench(self, device: int) -> PerfMetrics: + def matrix_f6f4_bench(self, device: int) -> float: return self.matrix_bench(device, "F6F4", "FLOP", "GFLOPS") - def matrix_mxf8_bench(self, device: int) -> PerfMetrics: + def matrix_mxf8_bench(self, device: int) -> float: return self.matrix_bench(device, "MXF8", "FLOP", "GFLOPS") - def fp16_benchmark(self, device: int) -> PerfMetrics: + def fp16_benchmark(self, device: int) -> float: return self.flops_bench(device, "FP16", "FLOP", "GFLOPS") - def bf16_benchmark(self, device: int) -> PerfMetrics: + def bf16_benchmark(self, device: int) -> float: return self.flops_bench(device, "BF16", "FLOP", "GFLOPS") - def fp32_benchmark(self, device: int) -> PerfMetrics: + def fp32_benchmark(self, device: int) -> float: return self.flops_bench(device, "FP32", "FLOP", "GFLOPS") - def fp64_benchmark(self, device: int) -> PerfMetrics: + def fp64_benchmark(self, device: int) -> float: return self.flops_bench(device, "FP64", "FLOP", "GFLOPS") - def int8_benchmark(self, device: int) -> PerfMetrics: + def int8_benchmark(self, device: int) -> float: return self.flops_bench(device, "INT8", "IOP", "GOPS") - def int32_benchmark(self, device: int) -> PerfMetrics: + def int32_benchmark(self, device: int) -> float: return self.flops_bench(device, "INT32", "IOP", "GOPS") - def int64_benchmark(self, device: int) -> PerfMetrics: + def int64_benchmark(self, device: int) -> float: return self.flops_bench(device, "INT64", "IOP", "GOPS") - def run_benchmark(self, device: int) -> dict[PerfMetrics]: + def run_benchmark(self, device: int) -> dict[str, float]: """ Run the roofline tests on the specified device. - Returns a dictionary of PerfMetrics. + Returns a dictionary mapping test name to its peak performance. """ with self.gpu_benchmark_lock(device): metrics_dict = {} @@ -821,16 +772,16 @@ def run_benchmark(self, device: int) -> dict[PerfMetrics]: for name, func in self.tests.items(): if name in self.unsupported_data_types: print(f"Skipping {name}") - metrics = PerfMetrics(0, 0, 0) + peak = 0.0 else: - metrics = func(device) + peak = func(device) - metrics_dict[name] = metrics + metrics_dict[name] = peak print("GPU Benchmarking completed") return metrics_dict - def dump_csv(self, metrics: dict[PerfMetrics], file_path: str) -> None: + def dump_csv(self, metrics: dict[str, float], file_path: str) -> None: """Generate a csv file containing the collected benchmark metrics.""" # TODO: Better way to map CSV column names? with open(file_path, "w", newline="", encoding="utf-8") as f: @@ -842,15 +793,11 @@ def dump_csv(self, metrics: dict[PerfMetrics], file_path: str) -> None: row = ["device"] for t in types: row.append(self.csv_cols_map[t]) - row.append(self.csv_cols_map[t] + "Low") - row.append(self.csv_cols_map[t] + "High") writer.writerow(row) row = [self.device_id] for t in types: - row.append(metrics[t].mean) - row.append(metrics[t].low) - row.append(metrics[t].high) + row.append(metrics[t]) writer.writerow(row) From 3a55d0b7448546a319745f6ed749d6dc7350ff90 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 16:50:00 +0100 Subject: [PATCH 23/45] Fixes and minor improvements --- .../docs/how-to/analyze/cli.rst | 2 +- .../benchmark/gfx9/benchmark_gfx950.py | 4 ++-- .../src/roofline/roofline_hover.py | 19 ++++++++++++------- .../src/roofline/roofline_html.py | 3 ++- .../src/roofline/roofline_main.py | 16 +++++++--------- .../tests/unit/roofline/test_roofline_main.py | 9 ++++++--- 6 files changed, 30 insertions(+), 23 deletions(-) diff --git a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst index 191e413fce94..46ec591e395b 100644 --- a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst +++ b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst @@ -573,7 +573,7 @@ Example with multiple ``--mem-level`` and ``--roofline-data-type`` options: Interactive Roofline HTML: -* Use the *AI axis* selector to choose one memory level per kernel, or *All peaks* to plot every level at once. Isolating one kernel shows it across all available memory levels. +* Use the *AI axis* selector to choose one memory level per kernel, or *All peaks* to plot every level at once. *Limited by* -- the default -- plots each kernel only at the roof that bounds it, so every point shows where that kernel is actually constrained. Isolating one kernel shows it across all available memory levels. * Use the *Precision* selector to choose one or more arithmetic precision's peak roofline to display on the plot. * Use the *Kernels* and *Bandwidth rooflines* panels to isolate, multi-select, or reset plotted items. The *Runtime shown* slider filters to the heaviest kernels that reach the selected GPU resident-time cutoff. * Hover over kernel dots and rooflines to see arithmetic intensity, throughput, roofline percentage, limiter, runtime, bandwidth, and compute-peak details. diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py b/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py index 49337c2d4cbb..61421d87d25f 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py @@ -100,12 +100,12 @@ def set_kernel_source(self) -> None: extern "C" __global__ void mfma_bf16(int iter, float *dummy) { vec16 result = {0}; - vec8<__fp16> a; + vec8 a; a[7] = a[6] = a[5] = a[4] = a[3] = a[2] = a[1] = a[0] = threadIdx.x; for(int i = 0; i < iter; ++i) { - result = __builtin_amdgcn_mfma_f32_32x32x16_f16(a, a, result, 0, 0, 0); + result = __builtin_amdgcn_mfma_f32_32x32x16_bf16(a, a, result, 0, 0, 0); } if (result[0] != 2*result[0]) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index b63f649c1520..b508b57a1c8e 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -40,8 +40,8 @@ def build_kernel_hover_template( name_html: str, limiter: str, limiter_category: str, - count: Optional[float], - total_dispatches: Optional[float], + kernel_dispatches_count: Optional[float], + total_dispatches_count: Optional[float], total_time: Optional[float], time_unit: str, pct_runtime: Optional[float], @@ -58,22 +58,27 @@ def build_kernel_hover_template( total_app_time_txt = format_hover_number(total_app_time, ",.0f") pct_runtime_txt = format_hover_number(pct_runtime, ",.2f") pct_dispatches = ( - 100.0 * count / total_dispatches - if count is not None and total_dispatches + 100.0 * kernel_dispatches_count / total_dispatches_count + if kernel_dispatches_count is not None and total_dispatches_count else None ) pct_dispatches_txt = format_hover_number(pct_dispatches, ",.2f") + count_txt = _format_integer(kernel_dispatches_count) + total_dispatches_txt = _format_integer(total_dispatches_count) return _hover( name_html, [ f"Limited by {limiter_category}: {limiter}", - f"Performance: %{{customdata[1]}}% (%{{y:,.0f}} / %{{customdata[0]}} {unit})", + f"Performance: %{{customdata[1]}}% " + f"(%{{y:,.0f}} / %{{customdata[0]}} {unit})", "AI: %{x:.6g}", "", "Details", "%{customdata[2]}", - f"Dispatch Count: {pct_dispatches_txt}% ({_format_integer(count)} / {_format_integer(total_dispatches)})", - f"Duration: {pct_runtime_txt}% ({kernel_time_txt} / {total_app_time_txt} {time_unit})", + f"Dispatch Count: {pct_dispatches_txt}% " + f"({count_txt} / {total_dispatches_txt})", + f"Duration: {pct_runtime_txt}% " + f"({kernel_time_txt} / {total_app_time_txt} {time_unit})", ], ) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_html.py b/projects/rocprofiler-compute/src/roofline/roofline_html.py index b0502c296566..7c8d70ee28d1 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_html.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_html.py @@ -107,7 +107,8 @@ def build_interactive_document( "Plot each kernel at its arithmetic intensity for this memory level, " "matching the (AI axis) marker in the Bandwidth rooflines panel. " "All peaks plots every level at once. Limited by plots only each " - "kernel's bottleneck point." + "kernel's bottleneck point; as with All peaks, no single roof is " + "marked as the AI axis." ), RUNTIME_TITLE=html.escape( "Show only the heaviest kernels whose combined percent of GPU " diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 1e6eb881de6c..4b255a0db555 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -381,7 +381,9 @@ def _determine_kernel_limiter( if not candidates: return "Unknown", "Unknown", None - roof_value, label, category = min(candidates, key=lambda candidate: candidate[0]) + roof_value, label, category = min( + candidates, key=lambda candidate: candidate[0] + ) return label, category, roof_value def _build_kernel_traces( @@ -424,9 +426,7 @@ def _build_kernel_traces( if limiter_category == "Memory" else min(points, key=lambda point: point["ai"])["peak"] ) - pct_roof = ( - 100.0 * points[0]["perf"] / roof_value if roof_value else None - ) + pct_roof = 100.0 * points[0]["perf"] / roof_value if roof_value else None hover_cells = [ format_hover_number(roof_value, ",.0f"), format_hover_number(pct_roof, ",.2f"), @@ -452,8 +452,8 @@ def _build_kernel_traces( name_html=wrap_hover_name(truncate_kernel_name(kernel_name)), limiter=limiter, limiter_category=limiter_category, - count=count_val, - total_dispatches=total_dispatches, + kernel_dispatches_count=count_val, + total_dispatches_count=total_dispatches, total_time=time_val, time_unit=time_unit, pct_runtime=pct_val, @@ -541,9 +541,7 @@ def _build_bandwidth_hover_html(bandwidth_entries: list[dict[str, Any]]) -> str: peak_bandwidth = entry["peak_bandwidth"] achieved_bandwidth = entry["achieved_bandwidth"] pct_bandwidth = ( - 100.0 * achieved_bandwidth / peak_bandwidth - if peak_bandwidth - else None + 100.0 * achieved_bandwidth / peak_bandwidth if peak_bandwidth else None ) lines.append( f"\u2003{entry['level_name']}: " diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 7caf8c6370da..bfe1f52cf023 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -181,7 +181,9 @@ def pct_roof(kernel: dict, point_index: int = 0) -> float: return float(kernel["points"][point_index]["hoverCells"][1]) -def test_kernel_traces_score_against_the_roof_that_binds_not_the_tallest_drawn() -> None: +def test_kernel_traces_score_against_the_roof_that_binds_not_the_tallest_drawn() -> ( + None +): """The percentage is scored against whichever roof the limiter actually names, not always the tallest compute roof stacked on the figure: dropping the tall MFMA peak from the candidate set changes both the label (falls @@ -744,8 +746,9 @@ def test_standalone_roofs_open_at_the_selected_cap(benchmarked_roofline) -> None assert roof["sampleAi"][-1] == pytest.approx(max(peaks) / roof["bandwidth"]) -def test_combined_figures_keep_every_ceiling(benchmarked_roofline) -> None: - """Narrowing the standalone document must not alter the source figures.""" +def test_dash_figures_keep_every_ceiling(benchmarked_roofline) -> None: + """The WebUI has no precision selector, so narrowing the standalone document + must not reach back into the figures Dash renders.""" roofline, flops_figure = stacked_figure(benchmarked_roofline, ["BF16", "FP64"]) source_model = roofline._Roofline__view_models["FLOP"] ceiling_indices = [trace["traceIndex"] for trace in source_model.compute_traces] From 953790d6307cddcc2fc9e4cde4738469dbe9f590 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 17:12:10 +0100 Subject: [PATCH 24/45] duplicated function is removed --- .../src/roofline/assets/roofline_plot.js | 10 +--------- .../rocprofiler-compute/src/roofline/roofline_main.py | 1 + 2 files changed, 2 insertions(+), 9 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 31a0480f952a..f3b3d3047825 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -24,7 +24,6 @@ // ---- Own presentation --------------------------------------------------- var ALL_PEAKS_LABEL = "All peaks"; var LIMITING_PEAK_LABEL = "Limited by"; - var KERNEL_LIST_NAME_MAX_LENGTH = 200; var FALLBACK_COLOR = "#888888"; var PLOT_DIM_OPACITY = 0.15; // How far a roof drawn at 45 degrees in data space may lean on screen. @@ -208,13 +207,6 @@ return "(" + shown + " / " + total + ")"; } - function truncateKernelListName(name) { - if (name.length <= KERNEL_LIST_NAME_MAX_LENGTH) { - return name; - } - return name.slice(0, KERNEL_LIST_NAME_MAX_LENGTH - 3) + "..."; - } - // A native `title` attribute gets dismissed by the row's click-to-filter // handler and won't reappear until the pointer leaves and re-enters, so the // full kernel name uses this mouse-following tooltip instead. @@ -1608,7 +1600,7 @@ kernelList.appendChild( createPanelRow({ color: kernel.color, - label: truncateKernelListName(kernel.name), + label: kernel.label, title: kernel.name, labelClass: "roofline-panel-name roofline-kernel-name", dataset: { index: String(index) }, diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 4b255a0db555..14276d9e3837 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -463,6 +463,7 @@ def _build_kernel_traces( ) kernels_model.append({ "name": kernel_name, + "label": truncate_kernel_name(kernel_name), "color": color, "points": points, "pctRuntime": pct_val, From 275e8f0275679b09a189457875c9680935f4a58e Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 17:22:19 +0100 Subject: [PATCH 25/45] tooltip implementation moved to a separate file --- .../assets/roofline_kernel_list_tooltip.css | 30 ++++++ .../assets/roofline_kernel_list_tooltip.js | 92 +++++++++++++++++++ .../src/roofline/assets/roofline_plot.css | 28 ------ .../src/roofline/assets/roofline_plot.js | 77 +--------------- .../src/roofline/roofline_html.py | 8 +- 5 files changed, 132 insertions(+), 103 deletions(-) create mode 100644 projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.css create mode 100644 projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.css b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.css new file mode 100644 index 000000000000..3bfed291bdb0 --- /dev/null +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.css @@ -0,0 +1,30 @@ +/* Copyright (c) Advanced Micro Devices, Inc. */ +/* SPDX-License-Identifier: MIT */ + +.roofline-name-tooltip { + position: fixed; + top: 0; + left: 0; + z-index: 500; + max-width: min(480px, calc(100vw - 24px)); + padding: 6px 10px; + background: var(--roofline-surface); + color: var(--roofline-text); + border: 1px solid var(--roofline-border); + border-radius: 6px; + box-shadow: 0 4px 12px rgba(0, 0, 0, 0.15); + font-family: var(--roofline-kernel-font, monospace); + font-size: 12px; + line-height: 1.35; + word-break: break-word; + overflow-wrap: anywhere; + pointer-events: none; + opacity: 0; + visibility: hidden; + transition: opacity 0.1s ease; +} + +.roofline-name-tooltip.visible { + opacity: 1; + visibility: visible; +} diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js new file mode 100644 index 000000000000..b6664cd4c1a1 --- /dev/null +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js @@ -0,0 +1,92 @@ +// Copyright (c) Advanced Micro Devices, Inc. +// SPDX-License-Identifier: MIT + +// A native `title` attribute gets dismissed by a kernel row's click-to-filter +// handler and won't reappear until the pointer leaves and re-enters, so the +// kernel list uses this mouse-following tooltip to show the full kernel name +// instead. +(function () { + "use strict"; + + var NAME_TOOLTIP_OFFSET = 14; + var NAME_TOOLTIP_DELAY_MS = 2000; + + var nameTooltip = null; + var nameTooltipTimer = null; + + function clamp(value, minimum, maximum) { + return Math.min(Math.max(value, minimum), maximum); + } + + function ensureNameTooltip() { + if (!nameTooltip) { + nameTooltip = document.createElement("div"); + nameTooltip.className = "roofline-name-tooltip"; + nameTooltip.setAttribute("role", "tooltip"); + document.body.appendChild(nameTooltip); + } + return nameTooltip; + } + + function positionNameTooltip(x, y) { + var tooltip = nameTooltip; + if (!tooltip) { + return; + } + var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); + var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); + tooltip.style.left = clamp(x + NAME_TOOLTIP_OFFSET, 4, maxLeft) + "px"; + tooltip.style.top = clamp(y + NAME_TOOLTIP_OFFSET, 4, maxTop) + "px"; + } + + function showNameTooltip(text, x, y) { + var tooltip = ensureNameTooltip(); + tooltip.textContent = text; + tooltip.classList.add("visible"); + positionNameTooltip(x, y); + } + + function scheduleNameTooltip(text, getPosition) { + clearTimeout(nameTooltipTimer); + nameTooltipTimer = setTimeout(function () { + var position = getPosition(); + showNameTooltip(text, position.x, position.y); + }, NAME_TOOLTIP_DELAY_MS); + } + + function hideNameTooltip() { + clearTimeout(nameTooltipTimer); + if (nameTooltip) { + nameTooltip.classList.remove("visible"); + } + } + + function attach(label, action, title) { + var lastMousePosition = { x: 0, y: 0 }; + label.addEventListener("mouseenter", function (event) { + lastMousePosition = { x: event.clientX, y: event.clientY }; + scheduleNameTooltip(title, function () { + return lastMousePosition; + }); + }); + label.addEventListener("mousemove", function (event) { + lastMousePosition = { x: event.clientX, y: event.clientY }; + if (nameTooltip && nameTooltip.classList.contains("visible")) { + positionNameTooltip(event.clientX, event.clientY); + } + }); + label.addEventListener("mouseleave", hideNameTooltip); + action.addEventListener("focus", function () { + var rect = label.getBoundingClientRect(); + scheduleNameTooltip(title, function () { + return { x: rect.left, y: rect.bottom }; + }); + }); + action.addEventListener("blur", hideNameTooltip); + } + + window.RooflineKernelListTooltip = { + attach: attach, + hide: hideNameTooltip, + }; +})(); diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css index 0aa6105afba0..583becfa808e 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.css @@ -499,34 +499,6 @@ body { cursor: pointer; } -.roofline-name-tooltip { - position: fixed; - top: 0; - left: 0; - z-index: 500; - max-width: min(480px, calc(100vw - 24px)); - padding: 6px 10px; - background: var(--roofline-surface); - color: var(--roofline-text); - border: 1px solid var(--roofline-border); - border-radius: 6px; - box-shadow: 0 4px 12px rgba(0, 0, 0, 0.15); - font-family: var(--roofline-kernel-font, monospace); - font-size: 12px; - line-height: 1.35; - word-break: break-word; - overflow-wrap: anywhere; - pointer-events: none; - opacity: 0; - visibility: hidden; - transition: opacity 0.1s ease; -} - -.roofline-name-tooltip.visible { - opacity: 1; - visibility: visible; -} - @media (max-width: 900px) { .roofline-body { flex-direction: column; diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index f3b3d3047825..4a2b0068acc5 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -207,57 +207,6 @@ return "(" + shown + " / " + total + ")"; } - // A native `title` attribute gets dismissed by the row's click-to-filter - // handler and won't reappear until the pointer leaves and re-enters, so the - // full kernel name uses this mouse-following tooltip instead. - var nameTooltip = null; - var nameTooltipTimer = null; - var NAME_TOOLTIP_OFFSET = 14; - var NAME_TOOLTIP_DELAY_MS = 2000; - - function ensureNameTooltip() { - if (!nameTooltip) { - nameTooltip = document.createElement("div"); - nameTooltip.className = "roofline-name-tooltip"; - nameTooltip.setAttribute("role", "tooltip"); - document.body.appendChild(nameTooltip); - } - return nameTooltip; - } - - function positionNameTooltip(x, y) { - var tooltip = nameTooltip; - if (!tooltip) { - return; - } - var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); - var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); - tooltip.style.left = clamp(x + NAME_TOOLTIP_OFFSET, 4, maxLeft) + "px"; - tooltip.style.top = clamp(y + NAME_TOOLTIP_OFFSET, 4, maxTop) + "px"; - } - - function showNameTooltip(text, x, y) { - var tooltip = ensureNameTooltip(); - tooltip.textContent = text; - tooltip.classList.add("visible"); - positionNameTooltip(x, y); - } - - function scheduleNameTooltip(text, getPosition) { - clearTimeout(nameTooltipTimer); - nameTooltipTimer = setTimeout(function () { - var position = getPosition(); - showNameTooltip(text, position.x, position.y); - }, NAME_TOOLTIP_DELAY_MS); - } - - function hideNameTooltip() { - clearTimeout(nameTooltipTimer); - if (nameTooltip) { - nameTooltip.classList.remove("visible"); - } - } - function eachKernelRow(fn) { if (!kernelList) { return; @@ -1471,27 +1420,7 @@ label.className = opts.labelClass || "roofline-panel-name"; label.textContent = opts.label; if (opts.title) { - var lastMousePosition = { x: 0, y: 0 }; - label.addEventListener("mouseenter", function (event) { - lastMousePosition = { x: event.clientX, y: event.clientY }; - scheduleNameTooltip(opts.title, function () { - return lastMousePosition; - }); - }); - label.addEventListener("mousemove", function (event) { - lastMousePosition = { x: event.clientX, y: event.clientY }; - if (nameTooltip && nameTooltip.classList.contains("visible")) { - positionNameTooltip(event.clientX, event.clientY); - } - }); - label.addEventListener("mouseleave", hideNameTooltip); - action.addEventListener("focus", function () { - var rect = label.getBoundingClientRect(); - scheduleNameTooltip(opts.title, function () { - return { x: rect.left, y: rect.bottom }; - }); - }); - action.addEventListener("blur", hideNameTooltip); + window.RooflineKernelListTooltip.attach(label, action, opts.title); } action.appendChild(swatch); @@ -1500,7 +1429,9 @@ action.appendChild(node); }); action.addEventListener("click", function (event) { - hideNameTooltip(); + if (opts.title) { + window.RooflineKernelListTooltip.hide(); + } opts.onClick(event); }); item.appendChild(action); diff --git a/projects/rocprofiler-compute/src/roofline/roofline_html.py b/projects/rocprofiler-compute/src/roofline/roofline_html.py index 7c8d70ee28d1..2278c092d1d5 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_html.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_html.py @@ -115,8 +115,12 @@ def build_interactive_document( "resident time reaches this cutoff. The rightmost stop shows every " "plotted kernel." ), - CSS=_read_asset("roofline_plot.css"), + CSS=_read_asset("roofline_plot.css") + + "\n" + + _read_asset("roofline_kernel_list_tooltip.css"), PLOT_FRAGMENT=fragment, MODEL_JSON=view_model.to_json(), - JS=_read_asset("roofline_plot.js"), + JS=_read_asset("roofline_kernel_list_tooltip.js") + + "\n" + + _read_asset("roofline_plot.js"), ) From b75b014df58c1c7a8bb3214152d9643470631c0f Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 17:29:37 +0100 Subject: [PATCH 26/45] minor renames --- .../assets/roofline_kernel_list_tooltip.js | 16 +++++++--------- .../src/roofline/assets/roofline_plot.js | 4 ++-- 2 files changed, 9 insertions(+), 11 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js index b6664cd4c1a1..c69651322121 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js @@ -1,15 +1,13 @@ // Copyright (c) Advanced Micro Devices, Inc. // SPDX-License-Identifier: MIT -// A native `title` attribute gets dismissed by a kernel row's click-to-filter -// handler and won't reappear until the pointer leaves and re-enters, so the -// kernel list uses this mouse-following tooltip to show the full kernel name -// instead. +// Kernel's list "click-to-filter" action prevents a default tooltip being shown. +// Therefore custom tooltip is added to overcome this. (function () { "use strict"; - var NAME_TOOLTIP_OFFSET = 14; - var NAME_TOOLTIP_DELAY_MS = 2000; + var nameTooltipOffset = 14; + var nameTooltipDelayMs = 2000; var nameTooltip = null; var nameTooltipTimer = null; @@ -35,8 +33,8 @@ } var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); - tooltip.style.left = clamp(x + NAME_TOOLTIP_OFFSET, 4, maxLeft) + "px"; - tooltip.style.top = clamp(y + NAME_TOOLTIP_OFFSET, 4, maxTop) + "px"; + tooltip.style.left = clamp(x + nameTooltipOffset, 4, maxLeft) + "px"; + tooltip.style.top = clamp(y + nameTooltipOffset, 4, maxTop) + "px"; } function showNameTooltip(text, x, y) { @@ -51,7 +49,7 @@ nameTooltipTimer = setTimeout(function () { var position = getPosition(); showNameTooltip(text, position.x, position.y); - }, NAME_TOOLTIP_DELAY_MS); + }, nameTooltipDelayMs); } function hideNameTooltip() { diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 4a2b0068acc5..609f011f7d42 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -1420,7 +1420,7 @@ label.className = opts.labelClass || "roofline-panel-name"; label.textContent = opts.label; if (opts.title) { - window.RooflineKernelListTooltip.attach(label, action, opts.title); + RooflineKernelListTooltip.attach(label, action, opts.title); } action.appendChild(swatch); @@ -1430,7 +1430,7 @@ }); action.addEventListener("click", function (event) { if (opts.title) { - window.RooflineKernelListTooltip.hide(); + RooflineKernelListTooltip.hide(); } opts.onClick(event); }); From 535cf8469b513ea8caf7b8ea75d7aa7f2401d78a Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 17:47:40 +0100 Subject: [PATCH 27/45] tooltip is a class --- .../docs/how-to/analyze/cli.rst | 3 +- .../assets/roofline_kernel_list_tooltip.js | 137 ++++++++++-------- 2 files changed, 76 insertions(+), 64 deletions(-) diff --git a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst index 46ec591e395b..1d6b8db2a4e1 100644 --- a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst +++ b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst @@ -573,7 +573,8 @@ Example with multiple ``--mem-level`` and ``--roofline-data-type`` options: Interactive Roofline HTML: -* Use the *AI axis* selector to choose one memory level per kernel, or *All peaks* to plot every level at once. *Limited by* -- the default -- plots each kernel only at the roof that bounds it, so every point shows where that kernel is actually constrained. Isolating one kernel shows it across all available memory levels. +* Use the *AI axis* selector to choose one memory level per kernel, *All peaks* to plot every level at once or *Limited by* to show levels which are bottleneck for each kernel. +Isolating one kernel shows it across all available memory levels. * Use the *Precision* selector to choose one or more arithmetic precision's peak roofline to display on the plot. * Use the *Kernels* and *Bandwidth rooflines* panels to isolate, multi-select, or reset plotted items. The *Runtime shown* slider filters to the heaviest kernels that reach the selected GPU resident-time cutoff. * Hover over kernel dots and rooflines to see arithmetic intensity, throughput, roofline percentage, limiter, runtime, bandwidth, and compute-peak details. diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js index c69651322121..f99962f8b3ab 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js @@ -6,85 +6,96 @@ (function () { "use strict"; - var nameTooltipOffset = 14; - var nameTooltipDelayMs = 2000; - - var nameTooltip = null; - var nameTooltipTimer = null; - function clamp(value, minimum, maximum) { return Math.min(Math.max(value, minimum), maximum); } - function ensureNameTooltip() { - if (!nameTooltip) { - nameTooltip = document.createElement("div"); - nameTooltip.className = "roofline-name-tooltip"; - nameTooltip.setAttribute("role", "tooltip"); - document.body.appendChild(nameTooltip); + class RooflineKernelNameTooltip { + constructor() { + this.offset = 14; + this.delayMs = 2000; + this.element = null; + this.timer = null; } - return nameTooltip; - } - function positionNameTooltip(x, y) { - var tooltip = nameTooltip; - if (!tooltip) { - return; + createElement() { + if (!this.element) { + this.element = document.createElement("div"); + this.element.className = "roofline-name-tooltip"; + this.element.setAttribute("role", "tooltip"); + document.body.appendChild(this.element); + } + return this.element; } - var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); - var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); - tooltip.style.left = clamp(x + nameTooltipOffset, 4, maxLeft) + "px"; - tooltip.style.top = clamp(y + nameTooltipOffset, 4, maxTop) + "px"; - } - function showNameTooltip(text, x, y) { - var tooltip = ensureNameTooltip(); - tooltip.textContent = text; - tooltip.classList.add("visible"); - positionNameTooltip(x, y); - } + position(x, y) { + var tooltip = this.element; + if (!tooltip) { + return; + } + var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); + var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); + tooltip.style.left = clamp(x + this.offset, 4, maxLeft) + "px"; + tooltip.style.top = clamp(y + this.offset, 4, maxTop) + "px"; + } - function scheduleNameTooltip(text, getPosition) { - clearTimeout(nameTooltipTimer); - nameTooltipTimer = setTimeout(function () { - var position = getPosition(); - showNameTooltip(text, position.x, position.y); - }, nameTooltipDelayMs); - } + show(text, x, y) { + var tooltip = this.createElement(); + tooltip.textContent = text; + tooltip.classList.add("visible"); + this.position(x, y); + } - function hideNameTooltip() { - clearTimeout(nameTooltipTimer); - if (nameTooltip) { - nameTooltip.classList.remove("visible"); + showWithDelay(text, getPosition) { + var self = this; + clearTimeout(this.timer); + this.timer = setTimeout(function () { + var position = getPosition(); + self.show(text, position.x, position.y); + }, this.delayMs); } - } - function attach(label, action, title) { - var lastMousePosition = { x: 0, y: 0 }; - label.addEventListener("mouseenter", function (event) { - lastMousePosition = { x: event.clientX, y: event.clientY }; - scheduleNameTooltip(title, function () { - return lastMousePosition; - }); - }); - label.addEventListener("mousemove", function (event) { - lastMousePosition = { x: event.clientX, y: event.clientY }; - if (nameTooltip && nameTooltip.classList.contains("visible")) { - positionNameTooltip(event.clientX, event.clientY); + hide() { + clearTimeout(this.timer); + if (this.element) { + this.element.classList.remove("visible"); } - }); - label.addEventListener("mouseleave", hideNameTooltip); - action.addEventListener("focus", function () { - var rect = label.getBoundingClientRect(); - scheduleNameTooltip(title, function () { - return { x: rect.left, y: rect.bottom }; + } + + attach(label, action, title) { + var self = this; + var lastMousePosition = { x: 0, y: 0 }; + label.addEventListener("mouseenter", function (event) { + lastMousePosition = { x: event.clientX, y: event.clientY }; + self.showWithDelay(title, function () { + return lastMousePosition; + }); + }); + label.addEventListener("mousemove", function (event) { + lastMousePosition = { x: event.clientX, y: event.clientY }; + if (self.element && self.element.classList.contains("visible")) { + self.position(event.clientX, event.clientY); + } }); - }); - action.addEventListener("blur", hideNameTooltip); + label.addEventListener("mouseleave", function () { + self.hide(); + }); + action.addEventListener("focus", function () { + var rect = label.getBoundingClientRect(); + self.showWithDelay(title, function () { + return { x: rect.left, y: rect.bottom }; + }); + }); + action.addEventListener("blur", function () { + self.hide(); + }); + } } + var tooltip = new RooflineKernelNameTooltip(); + window.RooflineKernelListTooltip = { - attach: attach, - hide: hideNameTooltip, + attach: tooltip.attach.bind(tooltip), + hide: tooltip.hide.bind(tooltip), }; })(); From eefcbc228bd044e4b0068d91d2b443626e030d15 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 17:51:07 +0100 Subject: [PATCH 28/45] functions refactoring --- .../assets/roofline_kernel_list_tooltip.js | 106 +++++++++--------- 1 file changed, 52 insertions(+), 54 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js index f99962f8b3ab..2ddd0fe518aa 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js @@ -11,70 +11,24 @@ } class RooflineKernelNameTooltip { - constructor() { - this.offset = 14; - this.delayMs = 2000; - this.element = null; - this.timer = null; - } - - createElement() { - if (!this.element) { - this.element = document.createElement("div"); - this.element.className = "roofline-name-tooltip"; - this.element.setAttribute("role", "tooltip"); - document.body.appendChild(this.element); - } - return this.element; - } - - position(x, y) { - var tooltip = this.element; - if (!tooltip) { - return; - } - var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); - var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); - tooltip.style.left = clamp(x + this.offset, 4, maxLeft) + "px"; - tooltip.style.top = clamp(y + this.offset, 4, maxTop) + "px"; - } - - show(text, x, y) { - var tooltip = this.createElement(); - tooltip.textContent = text; - tooltip.classList.add("visible"); - this.position(x, y); - } - - showWithDelay(text, getPosition) { - var self = this; - clearTimeout(this.timer); - this.timer = setTimeout(function () { - var position = getPosition(); - self.show(text, position.x, position.y); - }, this.delayMs); - } - - hide() { - clearTimeout(this.timer); - if (this.element) { - this.element.classList.remove("visible"); - } - } + #offset = 14; + #delayMs = 2000; + #element = null; + #timer = null; attach(label, action, title) { var self = this; var lastMousePosition = { x: 0, y: 0 }; label.addEventListener("mouseenter", function (event) { lastMousePosition = { x: event.clientX, y: event.clientY }; - self.showWithDelay(title, function () { + self.#showWithDelay(title, function () { return lastMousePosition; }); }); label.addEventListener("mousemove", function (event) { lastMousePosition = { x: event.clientX, y: event.clientY }; - if (self.element && self.element.classList.contains("visible")) { - self.position(event.clientX, event.clientY); + if (self.#element && self.#element.classList.contains("visible")) { + self.#position(event.clientX, event.clientY); } }); label.addEventListener("mouseleave", function () { @@ -82,7 +36,7 @@ }); action.addEventListener("focus", function () { var rect = label.getBoundingClientRect(); - self.showWithDelay(title, function () { + self.#showWithDelay(title, function () { return { x: rect.left, y: rect.bottom }; }); }); @@ -90,6 +44,50 @@ self.hide(); }); } + + hide() { + clearTimeout(this.#timer); + if (this.#element) { + this.#element.classList.remove("visible"); + } + } + + #showWithDelay(text, getPosition) { + var self = this; + clearTimeout(this.#timer); + this.#timer = setTimeout(function () { + var position = getPosition(); + self.#show(text, position.x, position.y); + }, this.#delayMs); + } + + #show(text, x, y) { + var tooltip = this.#createElement(); + tooltip.textContent = text; + tooltip.classList.add("visible"); + this.#position(x, y); + } + + #createElement() { + if (!this.#element) { + this.#element = document.createElement("div"); + this.#element.className = "roofline-name-tooltip"; + this.#element.setAttribute("role", "tooltip"); + document.body.appendChild(this.#element); + } + return this.#element; + } + + #position(x, y) { + var tooltip = this.#element; + if (!tooltip) { + return; + } + var maxLeft = Math.max(4, window.innerWidth - tooltip.offsetWidth - 4); + var maxTop = Math.max(4, window.innerHeight - tooltip.offsetHeight - 4); + tooltip.style.left = clamp(x + this.#offset, 4, maxLeft) + "px"; + tooltip.style.top = clamp(y + this.#offset, 4, maxTop) + "px"; + } } var tooltip = new RooflineKernelNameTooltip(); From a9a5788d2a4dd3601beead91c9ebb8611b9660b8 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 17:55:36 +0100 Subject: [PATCH 29/45] renames --- .../assets/roofline_kernel_list_tooltip.js | 20 +++++++++---------- .../src/roofline/assets/roofline_plot.js | 2 +- 2 files changed, 11 insertions(+), 11 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js index 2ddd0fe518aa..477568725855 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_kernel_list_tooltip.js @@ -16,31 +16,31 @@ #element = null; #timer = null; - attach(label, action, title) { + attachTooltipToKernelRow(kernelRow, kernelLabel, kernelFullName) { var self = this; var lastMousePosition = { x: 0, y: 0 }; - label.addEventListener("mouseenter", function (event) { + kernelLabel.addEventListener("mouseenter", function (event) { lastMousePosition = { x: event.clientX, y: event.clientY }; - self.#showWithDelay(title, function () { + self.#showWithDelay(kernelFullName, function () { return lastMousePosition; }); }); - label.addEventListener("mousemove", function (event) { + kernelLabel.addEventListener("mousemove", function (event) { lastMousePosition = { x: event.clientX, y: event.clientY }; if (self.#element && self.#element.classList.contains("visible")) { self.#position(event.clientX, event.clientY); } }); - label.addEventListener("mouseleave", function () { + kernelLabel.addEventListener("mouseleave", function () { self.hide(); }); - action.addEventListener("focus", function () { - var rect = label.getBoundingClientRect(); - self.#showWithDelay(title, function () { + kernelRow.addEventListener("focus", function () { + var rect = kernelLabel.getBoundingClientRect(); + self.#showWithDelay(kernelFullName, function () { return { x: rect.left, y: rect.bottom }; }); }); - action.addEventListener("blur", function () { + kernelRow.addEventListener("blur", function () { self.hide(); }); } @@ -93,7 +93,7 @@ var tooltip = new RooflineKernelNameTooltip(); window.RooflineKernelListTooltip = { - attach: tooltip.attach.bind(tooltip), + attachTooltipToKernelRow: tooltip.attachTooltipToKernelRow.bind(tooltip), hide: tooltip.hide.bind(tooltip), }; })(); diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 609f011f7d42..5540a91def89 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -1420,7 +1420,7 @@ label.className = opts.labelClass || "roofline-panel-name"; label.textContent = opts.label; if (opts.title) { - RooflineKernelListTooltip.attach(label, action, opts.title); + RooflineKernelListTooltip.attachTooltipToKernelRow(action, label, opts.title); } action.appendChild(swatch); From 149dd8081695e70056ac5af590c65916e21d2816 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 18:25:19 +0100 Subject: [PATCH 30/45] refactor --- .../src/roofline/roofline_hover.py | 65 +++++++++++-------- .../src/roofline/roofline_main.py | 13 ++-- 2 files changed, 47 insertions(+), 31 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index b508b57a1c8e..ba13f159758b 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -4,7 +4,7 @@ """Plotly hover-tooltip HTML for the roofline figure.""" import html -from typing import Optional +from typing import NamedTuple, Optional KERNEL_NAME_FONT_FAMILY = "ui-monospace, SFMono-Regular, Menlo, Consolas, monospace" @@ -36,35 +36,33 @@ def wrap_hover_name(name: str) -> str: ) +class DispatchCounts(NamedTuple): + """Per-kernel and total dispatch counts for the tooltip's count ratio.""" + + kernel: Optional[float] + total: Optional[float] + + +class DurationStats(NamedTuple): + """Per-kernel duration, its share of total runtime, and the time unit.""" + + total: Optional[float] + pct_runtime: Optional[float] + unit: str + + def build_kernel_hover_template( name_html: str, limiter: str, limiter_category: str, - kernel_dispatches_count: Optional[float], - total_dispatches_count: Optional[float], - total_time: Optional[float], - time_unit: str, - pct_runtime: Optional[float], + dispatches: DispatchCounts, + duration: DurationStats, ops_flops: str, ) -> str: """Kernel hover template; per-point values come from customdata.""" unit = f"G{ops_flops}s/s" - kernel_time_txt = format_hover_number(total_time, ",.0f") - total_app_time = ( - total_time / (pct_runtime / 100.0) - if total_time is not None and pct_runtime - else None - ) - total_app_time_txt = format_hover_number(total_app_time, ",.0f") - pct_runtime_txt = format_hover_number(pct_runtime, ",.2f") - pct_dispatches = ( - 100.0 * kernel_dispatches_count / total_dispatches_count - if kernel_dispatches_count is not None and total_dispatches_count - else None - ) - pct_dispatches_txt = format_hover_number(pct_dispatches, ",.2f") - count_txt = _format_integer(kernel_dispatches_count) - total_dispatches_txt = _format_integer(total_dispatches_count) + total_app_time = _invert_percent(duration.total, duration.pct_runtime) + pct_dispatches = _percent_of(dispatches.kernel, dispatches.total) return _hover( name_html, [ @@ -75,10 +73,11 @@ def build_kernel_hover_template( "", "Details", "%{customdata[2]}", - f"Dispatch Count: {pct_dispatches_txt}% " - f"({count_txt} / {total_dispatches_txt})", - f"Duration: {pct_runtime_txt}% " - f"({kernel_time_txt} / {total_app_time_txt} {time_unit})", + f"Dispatch Count: {format_hover_number(pct_dispatches, ',.2f')}% " + f"({_format_integer(dispatches.kernel)} / {_format_integer(dispatches.total)})", + f"Duration: {format_hover_number(duration.pct_runtime, ',.2f')}% " + f"({format_hover_number(duration.total, ',.0f')} / " + f"{format_hover_number(total_app_time, ',.0f')} {duration.unit})", ], ) @@ -139,6 +138,20 @@ def _format_integer(value: object) -> str: return f"{int(round(parsed)):,}" if parsed is not None else "N/A" +def _percent_of(part: Optional[float], whole: Optional[float]) -> Optional[float]: + """100 * part / whole, or None if not computable.""" + if part is None or not whole: + return None + return 100.0 * part / whole + + +def _invert_percent(value: Optional[float], pct: Optional[float]) -> Optional[float]: + """The whole that `value` is `pct` percent of, or None if not computable.""" + if value is None or not pct: + return None + return value * 100.0 / pct + + def _hover(header: str, rows: list[str]) -> str: """Assemble the common Plotly hover body.""" return "
".join([header, "", *rows]) + "" diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 14276d9e3837..9a5bba56e6a6 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -14,6 +14,8 @@ from roofline.roofline_frame import canonical_frame from roofline.roofline_hover import ( + DispatchCounts, + DurationStats, build_compute_peak_hover, build_kernel_hover_template, build_roof_hover, @@ -452,11 +454,12 @@ def _build_kernel_traces( name_html=wrap_hover_name(truncate_kernel_name(kernel_name)), limiter=limiter, limiter_category=limiter_category, - kernel_dispatches_count=count_val, - total_dispatches_count=total_dispatches, - total_time=time_val, - time_unit=time_unit, - pct_runtime=pct_val, + dispatches=DispatchCounts( + kernel=count_val, total=total_dispatches + ), + duration=DurationStats( + total=time_val, pct_runtime=pct_val, unit=time_unit + ), ops_flops=ops_flops, ), ) From be3a503a067d4e9b0411a1e8f227b0ac2a424cc9 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 18:33:13 +0100 Subject: [PATCH 31/45] rename --- .../src/roofline/roofline_hover.py | 14 +++++++------- .../src/roofline/roofline_main.py | 10 +++++----- 2 files changed, 12 insertions(+), 12 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index ba13f159758b..9e2b7a0a7a54 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -36,17 +36,17 @@ def wrap_hover_name(name: str) -> str: ) -class DispatchCounts(NamedTuple): +class KernelDispatchStats(NamedTuple): """Per-kernel and total dispatch counts for the tooltip's count ratio.""" kernel: Optional[float] total: Optional[float] -class DurationStats(NamedTuple): +class KernelDurationStats(NamedTuple): """Per-kernel duration, its share of total runtime, and the time unit.""" - total: Optional[float] + kernel: Optional[float] pct_runtime: Optional[float] unit: str @@ -55,13 +55,13 @@ def build_kernel_hover_template( name_html: str, limiter: str, limiter_category: str, - dispatches: DispatchCounts, - duration: DurationStats, + dispatches: KernelDispatchStats, + duration: KernelDurationStats, ops_flops: str, ) -> str: """Kernel hover template; per-point values come from customdata.""" unit = f"G{ops_flops}s/s" - total_app_time = _invert_percent(duration.total, duration.pct_runtime) + total_app_time = _invert_percent(duration.kernel, duration.pct_runtime) pct_dispatches = _percent_of(dispatches.kernel, dispatches.total) return _hover( name_html, @@ -76,7 +76,7 @@ def build_kernel_hover_template( f"Dispatch Count: {format_hover_number(pct_dispatches, ',.2f')}% " f"({_format_integer(dispatches.kernel)} / {_format_integer(dispatches.total)})", f"Duration: {format_hover_number(duration.pct_runtime, ',.2f')}% " - f"({format_hover_number(duration.total, ',.0f')} / " + f"({format_hover_number(duration.kernel, ',.0f')} / " f"{format_hover_number(total_app_time, ',.0f')} {duration.unit})", ], ) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 9a5bba56e6a6..6e310e526a8d 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -14,8 +14,8 @@ from roofline.roofline_frame import canonical_frame from roofline.roofline_hover import ( - DispatchCounts, - DurationStats, + KernelDispatchStats, + KernelDurationStats, build_compute_peak_hover, build_kernel_hover_template, build_roof_hover, @@ -454,11 +454,11 @@ def _build_kernel_traces( name_html=wrap_hover_name(truncate_kernel_name(kernel_name)), limiter=limiter, limiter_category=limiter_category, - dispatches=DispatchCounts( + dispatches=KernelDispatchStats( kernel=count_val, total=total_dispatches ), - duration=DurationStats( - total=time_val, pct_runtime=pct_val, unit=time_unit + duration=KernelDurationStats( + kernel=time_val, pct_runtime=pct_val, unit=time_unit ), ops_flops=ops_flops, ), From 0ff73237efe229d5397c1ce8cad13f48cf0ee08d Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 18:38:48 +0100 Subject: [PATCH 32/45] refactor --- .../src/roofline/roofline_hover.py | 19 ++++--------------- .../src/roofline/roofline_main.py | 9 ++++++++- 2 files changed, 12 insertions(+), 16 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index 9e2b7a0a7a54..84b1b7155c7c 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -37,17 +37,13 @@ def wrap_hover_name(name: str) -> str: class KernelDispatchStats(NamedTuple): - """Per-kernel and total dispatch counts for the tooltip's count ratio.""" - kernel: Optional[float] total: Optional[float] class KernelDurationStats(NamedTuple): - """Per-kernel duration, its share of total runtime, and the time unit.""" - kernel: Optional[float] - pct_runtime: Optional[float] + total: Optional[float] unit: str @@ -61,8 +57,8 @@ def build_kernel_hover_template( ) -> str: """Kernel hover template; per-point values come from customdata.""" unit = f"G{ops_flops}s/s" - total_app_time = _invert_percent(duration.kernel, duration.pct_runtime) pct_dispatches = _percent_of(dispatches.kernel, dispatches.total) + pct_runtime = _percent_of(duration.kernel, duration.total) return _hover( name_html, [ @@ -75,9 +71,9 @@ def build_kernel_hover_template( "%{customdata[2]}", f"Dispatch Count: {format_hover_number(pct_dispatches, ',.2f')}% " f"({_format_integer(dispatches.kernel)} / {_format_integer(dispatches.total)})", - f"Duration: {format_hover_number(duration.pct_runtime, ',.2f')}% " + f"Duration: {format_hover_number(pct_runtime, ',.2f')}% " f"({format_hover_number(duration.kernel, ',.0f')} / " - f"{format_hover_number(total_app_time, ',.0f')} {duration.unit})", + f"{format_hover_number(duration.total, ',.0f')} {duration.unit})", ], ) @@ -145,13 +141,6 @@ def _percent_of(part: Optional[float], whole: Optional[float]) -> Optional[float return 100.0 * part / whole -def _invert_percent(value: Optional[float], pct: Optional[float]) -> Optional[float]: - """The whole that `value` is `pct` percent of, or None if not computable.""" - if value is None or not pct: - return None - return value * 100.0 / pct - - def _hover(header: str, rows: list[str]) -> str: """Assemble the common Plotly hover body.""" return "
".join([header, "", *rows]) + "" diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 6e310e526a8d..0e334bf79ce5 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -141,6 +141,12 @@ def _decade_label(value: float) -> str: return f"1e{int(round(math.log10(value)))}" +def _invert_percent(value: Optional[float], pct: Optional[float]) -> Optional[float]: + if value is None or not pct: + return None + return value * 100.0 / pct + + def _frame_subtitle( bounds: tuple[float, float, float, float], is_machine_frame: bool ) -> str: @@ -420,6 +426,7 @@ def _build_kernel_traces( values[kernel_index] if kernel_index < len(values) else None for values in (kernel_colors, counts, total_time, pct_runtime) ) + total_app_time_val = _invert_percent(time_val, pct_val) limiter, limiter_category, roof_value = self._determine_kernel_limiter( level_ai, ceiling_data, points[0]["perf"], compute_peaks ) @@ -458,7 +465,7 @@ def _build_kernel_traces( kernel=count_val, total=total_dispatches ), duration=KernelDurationStats( - kernel=time_val, pct_runtime=pct_val, unit=time_unit + kernel=time_val, total=total_app_time_val, unit=time_unit ), ops_flops=ops_flops, ), From 716ad48a0abd230daf92e36278ece9979528f8a2 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 18:49:08 +0100 Subject: [PATCH 33/45] renames --- .../src/roofline/roofline_main.py | 16 +++++----------- .../src/utils/roofline_calc.py | 15 +++++++++------ .../tests/unit/utils/test_roofline_calc.py | 8 ++++---- 3 files changed, 18 insertions(+), 21 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 0e334bf79ce5..f6afa2c860e4 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -141,12 +141,6 @@ def _decade_label(value: float) -> str: return f"1e{int(round(math.log10(value)))}" -def _invert_percent(value: Optional[float], pct: Optional[float]) -> Optional[float]: - if value is None or not pct: - return None - return value * 100.0 / pct - - def _frame_subtitle( bounds: tuple[float, float, float, float], is_machine_frame: bool ) -> str: @@ -408,8 +402,9 @@ def _build_kernel_traces( kernels_model: list[dict[str, Any]] = [] counts = self.__ai_data.get("counts", []) - total_time = self.__ai_data.get("totalTime", []) - pct_runtime = self.__ai_data.get("pctRuntime", []) + kernel_time = self.__ai_data.get("kernelTime", []) + kernel_pct_runtime = self.__ai_data.get("kernelPctRuntime", []) + total_app_time = self.__ai_data.get("totalAppTime") time_unit = self.__ai_data.get("timeUnit", "") total_dispatches = sum(count for count in counts if count is not None) @@ -424,9 +419,8 @@ def _build_kernel_traces( color, count_val, time_val, pct_val = ( values[kernel_index] if kernel_index < len(values) else None - for values in (kernel_colors, counts, total_time, pct_runtime) + for values in (kernel_colors, counts, kernel_time, kernel_pct_runtime) ) - total_app_time_val = _invert_percent(time_val, pct_val) limiter, limiter_category, roof_value = self._determine_kernel_limiter( level_ai, ceiling_data, points[0]["perf"], compute_peaks ) @@ -465,7 +459,7 @@ def _build_kernel_traces( kernel=count_val, total=total_dispatches ), duration=KernelDurationStats( - kernel=time_val, total=total_app_time_val, unit=time_unit + kernel=time_val, total=total_app_time, unit=time_unit ), ops_flops=ops_flops, ), diff --git a/projects/rocprofiler-compute/src/utils/roofline_calc.py b/projects/rocprofiler-compute/src/utils/roofline_calc.py index 5f119c103da7..332ad90a0560 100644 --- a/projects/rocprofiler-compute/src/utils/roofline_calc.py +++ b/projects/rocprofiler-compute/src/utils/roofline_calc.py @@ -169,8 +169,9 @@ class PlotPoints: ai_lds: list[list[float]] kernelNames: list[str] counts: list[Optional[float]] - totalTime: list[Optional[float]] - pctRuntime: list[Optional[float]] + kernelTime: list[Optional[float]] + kernelPctRuntime: list[Optional[float]] + totalAppTime: Optional[float] timeUnit: str @classmethod @@ -184,8 +185,9 @@ def empty(cls) -> "PlotPoints": ai_lds=[[], []], kernelNames=[], counts=[], - totalTime=[], - pctRuntime=[], + kernelTime=[], + kernelPctRuntime=[], + totalAppTime=None, timeUnit="", ) @@ -532,8 +534,8 @@ def _append_kernel_point( getattr(plot_points, ai_field)[1].append(metrics["performance"]) plot_points.kernelNames.append(kernel_name) plot_points.counts.append(_stat_or_none(stat_row, "Count")) - plot_points.totalTime.append(_stat_or_none(stat_row, sum_column)) - plot_points.pctRuntime.append(_stat_or_none(stat_row, "Percent")) + plot_points.kernelTime.append(_stat_or_none(stat_row, sum_column)) + plot_points.kernelPctRuntime.append(_stat_or_none(stat_row, "Percent")) def calc_ai_analyze( @@ -566,6 +568,7 @@ def calc_ai_analyze( ) if sum_column: plot_points.timeUnit = sum_column[len("Sum(") : -1] + plot_points.totalAppTime = float(top_df[sum_column].sum()) kernel_ids_to_process = _resolve_kernel_ids(workload) console_debug("roofline", f"Found {len(kernel_ids_to_process)} kernels to process") diff --git a/projects/rocprofiler-compute/tests/unit/utils/test_roofline_calc.py b/projects/rocprofiler-compute/tests/unit/utils/test_roofline_calc.py index 5d59c62b62a9..67388fa10f88 100644 --- a/projects/rocprofiler-compute/tests/unit/utils/test_roofline_calc.py +++ b/projects/rocprofiler-compute/tests/unit/utils/test_roofline_calc.py @@ -231,14 +231,14 @@ def test_calc_ai_analyze_joins_per_kernel_stats( top_stats={"Count": 128, "Sum(ns)": 154000.0, "Percent": 12.4}, ) assert joined["counts"] == [128] - assert joined["totalTime"] == [154000.0] - assert joined["pctRuntime"] == [12.4] + assert joined["kernelTime"] == [154000.0] + assert joined["kernelPctRuntime"] == [12.4] assert joined["timeUnit"] == "ns", "the unit is only recoverable from the column" missing = run_calc_ai_analyze_with_values(monkeypatch, ai_values) assert missing["counts"] == [None] - assert missing["totalTime"] == [None] - assert missing["pctRuntime"] == [None] + assert missing["kernelTime"] == [None] + assert missing["kernelPctRuntime"] == [None] assert missing["timeUnit"] == "" From 740ae88cafa5a78ccdc9ca2f5626a339843112a1 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 18:53:44 +0100 Subject: [PATCH 34/45] refactor --- .../rocprofiler-compute/src/roofline/roofline_html.py | 5 ++--- .../rocprofiler-compute/src/roofline/roofline_main.py | 10 ---------- 2 files changed, 2 insertions(+), 13 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_html.py b/projects/rocprofiler-compute/src/roofline/roofline_html.py index 2278c092d1d5..82beeb8b8761 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_html.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_html.py @@ -106,9 +106,8 @@ def build_interactive_document( PEAK_TITLE=html.escape( "Plot each kernel at its arithmetic intensity for this memory level, " "matching the (AI axis) marker in the Bandwidth rooflines panel. " - "All peaks plots every level at once. Limited by plots only each " - "kernel's bottleneck point; as with All peaks, no single roof is " - "marked as the AI axis." + "'All peaks' plots every level at once." + "'Limited by' plots only each kernel's bottleneck points" ), RUNTIME_TITLE=html.escape( "Show only the heaviest kernels whose combined percent of GPU " diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index f6afa2c860e4..562c3c0e6e18 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -359,16 +359,6 @@ def _determine_kernel_limiter( performance: float, compute_peaks: list[tuple[str, float]], ) -> tuple[str, str, Optional[float]]: - """Name the specific binding roof for a kernel: the roof with the lowest - achievable performance at the kernel's operating point. Each stacked - datatype's compute peak, and each memory level's own bandwidth * AI, - is judged on its own, and only candidates the kernel's own measured - performance doesn't already exceed are eligible, so the roof named (and - its value) is one this kernel could actually be bound by, instead of an - already-impossible ceiling or always the tallest peak drawn. - - Returns (label, category, roof_value); category is "Memory" or - "Compute", roof_value is None when no candidate is eligible.""" candidates: list[tuple[float, str, str]] = [] for level_name, ai_value in level_ai.items(): bandwidth = self._peak_value(ceiling_data, level_name.lower()) From b21674407d91219c1d7b4766fd42fc96d103bbbe Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 18:58:49 +0100 Subject: [PATCH 35/45] refactor --- .../src/roofline/roofline_main.py | 74 +++++++++++-------- 1 file changed, 45 insertions(+), 29 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 562c3c0e6e18..a4d7bb61e00d 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -414,19 +414,8 @@ def _build_kernel_traces( limiter, limiter_category, roof_value = self._determine_kernel_limiter( level_ai, ceiling_data, points[0]["perf"], compute_peaks ) - limiting_peak = ( - limiter - if limiter_category == "Memory" - else min(points, key=lambda point: point["ai"])["peak"] - ) - pct_roof = 100.0 * points[0]["perf"] / roof_value if roof_value else None - hover_cells = [ - format_hover_number(roof_value, ",.0f"), - format_hover_number(pct_roof, ",.2f"), - bandwidth_html, - ] - for point in points: - point["hoverCells"] = hover_cells + limiting_peak = self._resolve_limiting_peak(limiter, limiter_category, points) + self._attach_kernel_hover_cells(points, roof_value, bandwidth_html) traces.append( go.Scatter( @@ -466,6 +455,29 @@ def _build_kernel_traces( return traces, kernels_model + @staticmethod + def _resolve_limiting_peak( + limiter: str, limiter_category: str, points: list[dict[str, Any]] + ) -> str: + if limiter_category == "Memory": + return limiter + return min(points, key=lambda point: point["ai"])["peak"] + + @staticmethod + def _attach_kernel_hover_cells( + points: list[dict[str, Any]], + roof_value: Optional[float], + bandwidth_html: str, + ) -> None: + pct_roof = 100.0 * points[0]["perf"] / roof_value if roof_value else None + hover_cells = [ + format_hover_number(roof_value, ",.0f"), + format_hover_number(pct_roof, ",.2f"), + bandwidth_html, + ] + for point in points: + point["hoverCells"] = hover_cells + def _build_kernel_points( self, kernel_index: int, @@ -494,35 +506,39 @@ def _build_kernel_points( performance = level_points[1][kernel_index] cache_key = cache_level.removeprefix("ai_") peak_bandwidth = self._peak_value(ceiling_data, cache_key) - - if not (ai_value > 0 and performance > 0): - # A level with no traffic still gets a 0-valued bandwidth - # line instead of silently disappearing from the tooltip. - bandwidth_entries.append({ - "level_name": level_name, - "achieved_bandwidth": 0.0, - "peak_bandwidth": peak_bandwidth, - }) - continue + has_traffic = ai_value > 0 and performance > 0 # This kernel's own achieved bandwidth at this level (not the # hardware ceiling): performance (FLOP/s) / AI (FLOP/Byte) = Byte/s. - achieved_bandwidth = performance / ai_value + # A level with no traffic still gets a 0-valued bandwidth line + # instead of silently disappearing from the tooltip. + achieved_bandwidth = performance / ai_value if has_traffic else 0.0 + bandwidth_entries.append( + self._build_bandwidth_entry(level_name, achieved_bandwidth, peak_bandwidth) + ) + if not has_traffic: + continue + points.append({ "peak": level_name, "ai": ai_value, "perf": performance, }) - bandwidth_entries.append({ - "level_name": level_name, - "achieved_bandwidth": achieved_bandwidth, - "peak_bandwidth": peak_bandwidth, - }) level_ai[level_name] = ai_value bandwidth_html = self._build_bandwidth_hover_html(bandwidth_entries) return points, level_ai, bandwidth_html + @staticmethod + def _build_bandwidth_entry( + level_name: str, achieved_bandwidth: float, peak_bandwidth: Optional[float] + ) -> dict[str, Any]: + return { + "level_name": level_name, + "achieved_bandwidth": achieved_bandwidth, + "peak_bandwidth": peak_bandwidth, + } + @staticmethod def _build_bandwidth_hover_html(bandwidth_entries: list[dict[str, Any]]) -> str: """One 'Bandwidth:' block listing every cache level this kernel has From 904b9bcaf90568eab6e9bf6e16cf6f24ee9b9813 Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 19:04:28 +0100 Subject: [PATCH 36/45] test refactoring --- .../tests/unit/roofline/test_roofline_main.py | 58 ++++++++++--------- 1 file changed, 32 insertions(+), 26 deletions(-) diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index bfe1f52cf023..6f2f331617fb 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -204,10 +204,9 @@ def test_kernel_traces_score_against_the_roof_that_binds_not_the_tallest_drawn() assert "Limited by Memory: HBM" in valu_traces[0].hovertemplate -def test_kernel_traces_name_the_specific_peak_not_the_tallest_stacked_one() -> None: - """When several compute peaks are stacked, the limiter names whichever one - actually bounds this kernel's own performance, not whichever peak happens - to be tallest on the figure.""" +def test_kernel_traces_given_multiple_stacked_compute_peaks__limiter_names_the_one_that_binds() -> ( + None +): ai_data = {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]} traces, _ = kernel_traces( @@ -219,13 +218,9 @@ def test_kernel_traces_name_the_specific_peak_not_the_tallest_stacked_one() -> N assert "Limited by Compute: FP32 MFMA" in traces[0].hovertemplate -def test_kernel_traces_ignore_a_memory_roof_the_kernel_already_exceeds() -> None: - """A memory level's implied ceiling (bandwidth * AI) that the kernel's own - measured performance already exceeds is not an achievable roof, so the - limiter must not name it -- exactly as compute peaks below performance are - already excluded. Here HBM implies a 1,500 GFLOP/s ceiling but the kernel - measures 5,000 GFLOP/s, so the limiter must fall through to the cheapest - eligible compute peak instead of naming the impossible memory ceiling.""" +def test_kernel_traces_given_memory_roof_already_exceeded__falls_back_to_compute_peak() -> ( + None +): traces, _ = kernel_traces( make_roofline(["FP32"]), {"ai_hbm": [[1.0], [5000.0]], "kernelNames": ["kA"]}, @@ -322,13 +317,9 @@ def test_kernel_traces_expose_the_leftmost_peak_when_the_limiter_is_unknown() -> assert model[0]["limitingPeak"] == "L2" -def test_kernel_traces_share_one_performance_value_across_a_kernels_points() -> None: - """A kernel's HBM point and L2 point measure the same achieved performance - but sit at different arithmetic intensities, so today each point is scored - against its own level's bandwidth -- 25.00% at HBM, 66.67% at L2. Once the - limiter's own closest-roof search (L2's 1,500 GFLOP/s implied ceiling beats - HBM's 4,000) drives the displayed percentage too, both of the kernel's - points must show that same 66.67% / 1,500, not their own per-level value.""" +def test_kernel_traces_given_differing_level_bandwidths__all_points_share_the_binding_roofs_percentage() -> ( + None +): ceiling = { "hbm": [[0.01, 1.0], [1.0, 2000.0], 2000.0], "l2": [[0.01, 1.0], [1.0, 3000.0], 3000.0], @@ -351,9 +342,7 @@ def test_kernel_traces_share_one_performance_value_across_a_kernels_points() -> assert points["L2"]["hoverCells"][:2] == ["1,500", "66.67"] -def test_bandwidth_hover_shows_zero_levels() -> None: - """Every cache level still shows up at 0 when a kernel does no traffic - there, instead of vanishing from the tooltip.""" +def test_bandwidth_hover_given_a_level_with_no_traffic__still_lists_it_at_zero() -> None: ceiling = dict(CEILING, lds=[[0.01, 1.0], [1.0, 800.0], 800.0]) _, model = kernel_traces( make_roofline(["FP32"]), @@ -374,11 +363,7 @@ def test_bandwidth_hover_shows_zero_levels() -> None: ) -def test_kernel_hover_carries_the_whole_name_up_to_the_limit() -> None: - """A demangled name within the tooltip's length cap reaches the tooltip - whole. It is wrapped onto as many lines as it takes, but nothing is - dropped: two instantiations of the same function are told apart by - template arguments that run to the very end of the name.""" +def test_kernel_hover_given_a_name_under_the_length_cap__wraps_it_whole() -> None: name = "Cijk_Alik_Bljk_" + "SB_MT256x256x16_MI32x32x2x1_" * 6 traces, _ = kernel_traces( @@ -760,6 +745,27 @@ def test_dash_figures_keep_every_ceiling(benchmarked_roofline) -> None: assert drawn_roof_knees(flops_figure) == roof_extents +def test_default_precisions_given_fp32_fp16_fp8_present__orders_fp32_before_fp16_before_fp8() -> ( + None +): + from roofline.roofline_main import _default_precisions + + assert _default_precisions(["FP16", "FP32", "FP8", "FP64"]) == [ + "FP32", + "FP16", + "FP8", + ] + assert _default_precisions(["FP8"]) == ["FP8"] + + +def test_default_precisions_given_none_preferred_present__falls_back_to_first_datatype() -> ( + None +): + from roofline.roofline_main import _default_precisions + + assert _default_precisions(["FP64", "BF16"]) == ["FP64"] + + def test_construct_plotly_figures_all_datatypes_ignores_cli_selection( benchmarked_roofline, ) -> None: From 1819e134400ec97e091a80547db673d5f691f0ba Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 19:48:20 +0100 Subject: [PATCH 37/45] test refactor --- .../tests/unit/roofline/test_roofline_main.py | 24 ------------------- 1 file changed, 24 deletions(-) diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 6f2f331617fb..938debb37526 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -181,29 +181,6 @@ def pct_roof(kernel: dict, point_index: int = 0) -> float: return float(kernel["points"][point_index]["hoverCells"][1]) -def test_kernel_traces_score_against_the_roof_that_binds_not_the_tallest_drawn() -> ( - None -): - """The percentage is scored against whichever roof the limiter actually - names, not always the tallest compute roof stacked on the figure: dropping - the tall MFMA peak from the candidate set changes both the label (falls - back to the memory level, since VALU's peak is below this kernel's own - performance) and the percentage (scored against HBM's own 150,000 - GFLOP/s implied ceiling instead of MFMA's 90,000), and neither case's - percentage exceeds 100%.""" - ai_data = {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]} - - matrix_traces, matrix_capped = kernel_traces(make_roofline(["FP32"]), ai_data) - valu_traces, valu_capped = kernel_traces( - make_roofline(["FP32"]), ai_data, compute_peaks=[("FP32 VALU", 9000.0)] - ) - - assert pct_roof(matrix_capped[0]) == pytest.approx(55.56, abs=0.01) - assert pct_roof(valu_capped[0]) == pytest.approx(33.33, abs=0.01) - assert "Limited by Compute: FP32 MFMA" in matrix_traces[0].hovertemplate - assert "Limited by Memory: HBM" in valu_traces[0].hovertemplate - - def test_kernel_traces_given_multiple_stacked_compute_peaks__limiter_names_the_one_that_binds() -> ( None ): @@ -227,7 +204,6 @@ def test_kernel_traces_given_memory_roof_already_exceeded__falls_back_to_compute ) assert "Limited by Compute: FP32 VALU" in traces[0].hovertemplate - assert "Limited by Memory: HBM" not in traces[0].hovertemplate def test_kernel_traces_name_the_roof_that_binds() -> None: From 723a8699d61143e1af9d32e5b58d8aa03d757e5a Mon Sep 17 00:00:00 2001 From: svolkov-amd Date: Thu, 24 Sep 2026 20:21:25 +0100 Subject: [PATCH 38/45] test refactor --- .../tests/unit/roofline/test_roofline_main.py | 19 ++++++++++--------- 1 file changed, 10 insertions(+), 9 deletions(-) diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 938debb37526..50e30121feef 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -238,10 +238,6 @@ def test_kernel_traces_name_the_roof_that_binds() -> None: def test_kernel_traces_expose_the_limiting_peak_for_a_memory_bound_kernel() -> None: - """A memory-bound kernel's model carries a structured limitingPeak field, - naming the same cache level the hover text already calls out as the - binding roof, so the client can pick that one point without - re-deriving the limiter itself.""" traces, model = kernel_traces( make_roofline(["FP32"]), {"ai_hbm": [[1.0], [900.0]], "kernelNames": ["kA"]}, @@ -332,11 +328,16 @@ def test_bandwidth_hover_given_a_level_with_no_traffic__still_lists_it_at_zero() ceiling_data=ceiling, ) bandwidth_html = model[0]["points"][0]["hoverCells"][2] - assert bandwidth_html == ( - "Bandwidth:
\u2003L2: N/A% (0.000 / N/A GB/s)" - "
\u2003HBM: 60.00% (0.900 / 1.500 TB/s)" - "
\u2003LDS: 0.00% (0.000 / 800.000 GB/s)" - ) + lines = bandwidth_html.split("
") + + lds_line = next(line for line in lines if "LDS" in line) + assert "0.000" in lds_line + + l2_line = next(line for line in lines if "L2" in line) + assert "0.000" in l2_line + + hbm_line = next(line for line in lines if "HBM" in line) + assert "900" in hbm_line def test_kernel_hover_given_a_name_under_the_length_cap__wraps_it_whole() -> None: From 5ad6ec9b5001f7ccf61d5dc57f9bce5fe5055df2 Mon Sep 17 00:00:00 2001 From: Carrie Fallows Date: Thu, 8 Oct 2026 14:10:31 -0400 Subject: [PATCH 39/45] Fix formatting- break up long lines, relabel test cases due to function name length exceeding limit Signed-off-by: Carrie Fallows --- .../src/roofline/roofline_hover.py | 3 ++- .../src/roofline/roofline_main.py | 8 +++++-- .../tests/unit/roofline/test_roofline_main.py | 22 +++++-------------- 3 files changed, 14 insertions(+), 19 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_hover.py b/projects/rocprofiler-compute/src/roofline/roofline_hover.py index 84b1b7155c7c..43e038cdb80d 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_hover.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_hover.py @@ -70,7 +70,8 @@ def build_kernel_hover_template( "Details", "%{customdata[2]}", f"Dispatch Count: {format_hover_number(pct_dispatches, ',.2f')}% " - f"({_format_integer(dispatches.kernel)} / {_format_integer(dispatches.total)})", + f"({_format_integer(dispatches.kernel)}" + f" / {_format_integer(dispatches.total)})", f"Duration: {format_hover_number(pct_runtime, ',.2f')}% " f"({format_hover_number(duration.kernel, ',.0f')} / " f"{format_hover_number(duration.total, ',.0f')} {duration.unit})", diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index a4d7bb61e00d..9ae2c5e9b0df 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -414,7 +414,9 @@ def _build_kernel_traces( limiter, limiter_category, roof_value = self._determine_kernel_limiter( level_ai, ceiling_data, points[0]["perf"], compute_peaks ) - limiting_peak = self._resolve_limiting_peak(limiter, limiter_category, points) + limiting_peak = self._resolve_limiting_peak( + limiter, limiter_category, points + ) self._attach_kernel_hover_cells(points, roof_value, bandwidth_html) traces.append( @@ -514,7 +516,9 @@ def _build_kernel_points( # instead of silently disappearing from the tooltip. achieved_bandwidth = performance / ai_value if has_traffic else 0.0 bandwidth_entries.append( - self._build_bandwidth_entry(level_name, achieved_bandwidth, peak_bandwidth) + self._build_bandwidth_entry( + level_name, achieved_bandwidth, peak_bandwidth + ) ) if not has_traffic: continue diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index 50e30121feef..c5f4ebaaf0a6 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -181,9 +181,7 @@ def pct_roof(kernel: dict, point_index: int = 0) -> float: return float(kernel["points"][point_index]["hoverCells"][1]) -def test_kernel_traces_given_multiple_stacked_compute_peaks__limiter_names_the_one_that_binds() -> ( - None -): +def test_kernel_traces_stacked_compute_peaks__limiter_names_binding_one() -> None: ai_data = {"ai_hbm": [[100.0], [50000.0]], "kernelNames": ["kA"]} traces, _ = kernel_traces( @@ -195,9 +193,7 @@ def test_kernel_traces_given_multiple_stacked_compute_peaks__limiter_names_the_o assert "Limited by Compute: FP32 MFMA" in traces[0].hovertemplate -def test_kernel_traces_given_memory_roof_already_exceeded__falls_back_to_compute_peak() -> ( - None -): +def test_kernel_traces_exceeded_memory_roof__falls_back_to_compute_peak() -> None: traces, _ = kernel_traces( make_roofline(["FP32"]), {"ai_hbm": [[1.0], [5000.0]], "kernelNames": ["kA"]}, @@ -289,9 +285,7 @@ def test_kernel_traces_expose_the_leftmost_peak_when_the_limiter_is_unknown() -> assert model[0]["limitingPeak"] == "L2" -def test_kernel_traces_given_differing_level_bandwidths__all_points_share_the_binding_roofs_percentage() -> ( - None -): +def test_kernel_traces_differing_bandwidths__points_share_binding_roof_pct() -> None: ceiling = { "hbm": [[0.01, 1.0], [1.0, 2000.0], 2000.0], "l2": [[0.01, 1.0], [1.0, 3000.0], 3000.0], @@ -314,7 +308,7 @@ def test_kernel_traces_given_differing_level_bandwidths__all_points_share_the_bi assert points["L2"]["hoverCells"][:2] == ["1,500", "66.67"] -def test_bandwidth_hover_given_a_level_with_no_traffic__still_lists_it_at_zero() -> None: +def test_bandwidth_hover_no_traffic_level__still_lists_at_zero() -> None: ceiling = dict(CEILING, lds=[[0.01, 1.0], [1.0, 800.0], 800.0]) _, model = kernel_traces( make_roofline(["FP32"]), @@ -722,9 +716,7 @@ def test_dash_figures_keep_every_ceiling(benchmarked_roofline) -> None: assert drawn_roof_knees(flops_figure) == roof_extents -def test_default_precisions_given_fp32_fp16_fp8_present__orders_fp32_before_fp16_before_fp8() -> ( - None -): +def test_default_precisions_fp32_fp16_fp8_present__orders_fp32_first() -> None: from roofline.roofline_main import _default_precisions assert _default_precisions(["FP16", "FP32", "FP8", "FP64"]) == [ @@ -735,9 +727,7 @@ def test_default_precisions_given_fp32_fp16_fp8_present__orders_fp32_before_fp16 assert _default_precisions(["FP8"]) == ["FP8"] -def test_default_precisions_given_none_preferred_present__falls_back_to_first_datatype() -> ( - None -): +def test_default_precisions_no_preferred__falls_back_to_first_datatype() -> None: from roofline.roofline_main import _default_precisions assert _default_precisions(["FP64", "BF16"]) == ["FP64"] From 98d10dce098bd6961edda4a8469389ffc6ad353d Mon Sep 17 00:00:00 2001 From: Carrie Fallows Date: Wed, 14 Oct 2026 01:33:53 -0400 Subject: [PATCH 40/45] =?UTF-8?q?roofline=5Fplot.js:=20Kernel=20sidebar=20?= =?UTF-8?q?buttons=20now=20get=20aria-label=20set=20to=20the=20full=20kern?= =?UTF-8?q?el=20name=20when=20truncated,=20so=20screen=20readers=20announc?= =?UTF-8?q?e=20the=20complete=20name.=20roofline=5Fcalc.py,=20roofline=5Fm?= =?UTF-8?q?ain.py:=20totalAppDispatches=20added=20to=20PlotPoints,=20popul?= =?UTF-8?q?ated=20from=20top=5Fdf["Count"].sum()=20(the=20full=20unfiltere?= =?UTF-8?q?d=20table).=20roofline=5Fmain.py=20now=20reads=20it=20from=20ai?= =?UTF-8?q?=5Fdata=20instead=20of=20summing=20the=20filtered=20counts=20li?= =?UTF-8?q?st=20=E2=80=94=20fixes=20the=20case=20where=20a=20selected=20ke?= =?UTF-8?q?rnel=20could=20show=20100%=20dispatch=20share.=20=20roofline=5F?= =?UTF-8?q?html.py,=20cli.rst:=20Missing=20space=20between=20adjacent=20st?= =?UTF-8?q?ring=20literals=20fixe,=20Inaccurate=20"bottleneck=20points"=20?= =?UTF-8?q?description=20replaced=20with=20accurate=20fallback=20wording?= =?UTF-8?q?=20in=20both=20the=20in-HTML=20tooltip=20and=20the=20docs.?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Signed-off-by: Carrie Fallows --- projects/rocprofiler-compute/docs/how-to/analyze/cli.rst | 2 +- .../rocprofiler-compute/src/roofline/assets/roofline_plot.js | 3 +++ projects/rocprofiler-compute/src/roofline/roofline_html.py | 5 +++-- projects/rocprofiler-compute/src/roofline/roofline_main.py | 2 +- projects/rocprofiler-compute/src/utils/roofline_calc.py | 4 ++++ 5 files changed, 12 insertions(+), 4 deletions(-) diff --git a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst index db1c0e1ba9e2..6ef8fe06170e 100644 --- a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst +++ b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst @@ -624,7 +624,7 @@ Example with multiple ``--mem-level`` and ``--roofline-data-type`` options: Interactive Roofline HTML: -* Use the *AI axis* selector to choose one memory level per kernel, *All peaks* to plot every level at once or *Limited by* to show levels which are bottleneck for each kernel. +* Use the *AI axis* selector to choose one memory level per kernel, *All peaks* to plot every level at once, or *Limited by* to show each kernel's limiting memory-level point (falling back to its leftmost point when no memory-level limiter is available). Isolating one kernel shows it across all available memory levels. * Use the *Precision* selector to choose one or more arithmetic precision's peak roofline to display on the plot. * Use the *Kernels* and *Bandwidth rooflines* panels to isolate, multi-select, or reset plotted items. The *Runtime shown* slider filters to the heaviest kernels that reach the selected GPU resident-time cutoff. diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 5540a91def89..0cc17ddbd83d 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -1411,6 +1411,9 @@ var action = document.createElement("button"); action.type = "button"; action.className = "roofline-panel-action"; + if (opts.title) { + action.setAttribute("aria-label", opts.title); + } var swatch = document.createElement("span"); swatch.className = opts.swatchClass || "roofline-swatch"; diff --git a/projects/rocprofiler-compute/src/roofline/roofline_html.py b/projects/rocprofiler-compute/src/roofline/roofline_html.py index 82beeb8b8761..145b3156d2c6 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_html.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_html.py @@ -106,8 +106,9 @@ def build_interactive_document( PEAK_TITLE=html.escape( "Plot each kernel at its arithmetic intensity for this memory level, " "matching the (AI axis) marker in the Bandwidth rooflines panel. " - "'All peaks' plots every level at once." - "'Limited by' plots only each kernel's bottleneck points" + "'All peaks' plots every level at once. " + "'Limited by' plots the limiting memory-level point, or the " + "leftmost point when no memory-level limiter is available." ), RUNTIME_TITLE=html.escape( "Show only the heaviest kernels whose combined percent of GPU " diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index 9ae2c5e9b0df..fa3b753740fe 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -395,8 +395,8 @@ def _build_kernel_traces( kernel_time = self.__ai_data.get("kernelTime", []) kernel_pct_runtime = self.__ai_data.get("kernelPctRuntime", []) total_app_time = self.__ai_data.get("totalAppTime") + total_dispatches = self.__ai_data.get("totalAppDispatches") time_unit = self.__ai_data.get("timeUnit", "") - total_dispatches = sum(count for count in counts if count is not None) for kernel_index, kernel_name in enumerate(kernel_names): points, level_ai, bandwidth_html = self._build_kernel_points( diff --git a/projects/rocprofiler-compute/src/utils/roofline_calc.py b/projects/rocprofiler-compute/src/utils/roofline_calc.py index 332ad90a0560..cca2cb7f6318 100644 --- a/projects/rocprofiler-compute/src/utils/roofline_calc.py +++ b/projects/rocprofiler-compute/src/utils/roofline_calc.py @@ -172,6 +172,7 @@ class PlotPoints: kernelTime: list[Optional[float]] kernelPctRuntime: list[Optional[float]] totalAppTime: Optional[float] + totalAppDispatches: Optional[float] timeUnit: str @classmethod @@ -188,6 +189,7 @@ def empty(cls) -> "PlotPoints": kernelTime=[], kernelPctRuntime=[], totalAppTime=None, + totalAppDispatches=None, timeUnit="", ) @@ -569,6 +571,8 @@ def calc_ai_analyze( if sum_column: plot_points.timeUnit = sum_column[len("Sum(") : -1] plot_points.totalAppTime = float(top_df[sum_column].sum()) + if "Count" in top_df.columns: + plot_points.totalAppDispatches = float(top_df["Count"].sum()) kernel_ids_to_process = _resolve_kernel_ids(workload) console_debug("roofline", f"Found {len(kernel_ids_to_process)} kernels to process") From dc9177673fd1dc2912bdfcb7bc5fb8a5cc678857 Mon Sep 17 00:00:00 2001 From: Carrie Fallows Date: Wed, 14 Oct 2026 12:36:17 -0400 Subject: [PATCH 41/45] Change BF16 MFMA builtin to match FP16 change. Create CHANGELOG entries. Signed-off-by: Carrie Fallows --- projects/rocprofiler-compute/CHANGELOG.md | 6 ++++++ .../src/roofline/benchmark/gfx9/benchmark_gfx950.py | 6 +++--- 2 files changed, 9 insertions(+), 3 deletions(-) diff --git a/projects/rocprofiler-compute/CHANGELOG.md b/projects/rocprofiler-compute/CHANGELOG.md index 82b1f8b12246..460edb3a2dfb 100644 --- a/projects/rocprofiler-compute/CHANGELOG.md +++ b/projects/rocprofiler-compute/CHANGELOG.md @@ -7,6 +7,8 @@ Full documentation for ROCm Compute Profiler is available at [https://rocm.docs. ### Added +* Added a *Limited by* mode to the *AI axis* selector in the standalone roofline HTML. It shows each kernel at its limiting memory-level point, falling back to the leftmost point for compute-bound or unknown kernels. + * Added Data Fabric to MALL bandwidth arrows to the CDNA (gfx9) Memory Chart. On gfx940–gfx942 these report estimated HBM bandwidth, derived from a 64B-per-request approximation, through the new `Estimated HBM Read BW` and `Estimated HBM Write and Atomic BW` metrics. On gfx950 they use the existing `HBM Read BW` and the new `HBM Write and Atomic BW` metric. * Added the `VL1 Coalesce` metric to the VL1D panel of the CDNA (gfx9) Memory Chart. @@ -15,6 +17,10 @@ Full documentation for ROCm Compute Profiler is available at [https://rocm.docs. ### Changed +* Improved the standalone roofline HTML kernel hover tooltip: it now shows the limiting roof, percent of roofline, achieved and peak bandwidth per memory level, dispatch count and share of total application dispatches, and kernel duration and share of total application time. Kernel names longer than 200 characters are truncated in the sidebar, with the full name shown on hover. + +* The roofline benchmark now reports the peak measured sample across runs instead of the mean. + * Renamed the Memory Chart left-hand panel from "Kernel" to "Compute Units" on CDNA (gfx9) and to "WGPs" on gfx115x and gfx1250. The panel now shows resource allocation stats: Wave Occupancy as a percent of the maximum wavefronts per CU (gfx9) or per WGP (gfx115x, gfx1250), vGPRs, sGPRs, Scratch in KB per wave, LDS Allocation, and Workgroups per active CU (gfx9) or per WGP (gfx115x, gfx1250). * All Memory Chart bandwidth values now use uniform fixed-point GB/s formatting (3 decimal places) for easy cross-level comparison. diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py b/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py index 61421d87d25f..c9821ebc1980 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/gfx9/benchmark_gfx950.py @@ -29,7 +29,7 @@ def __init__(self, device_id: int, cache_sizes: dict) -> None: "F8": 32768, "F16": 16384, "F32": 4096, - "BF16": 32768, + "BF16": 16384, "I8": 65536, "F64": 2048, } @@ -99,13 +99,13 @@ def set_kernel_source(self) -> None: + """ extern "C" __global__ void mfma_bf16(int iter, float *dummy) { - vec16 result = {0}; + vec4 result = {0}; vec8 a; a[7] = a[6] = a[5] = a[4] = a[3] = a[2] = a[1] = a[0] = threadIdx.x; for(int i = 0; i < iter; ++i) { - result = __builtin_amdgcn_mfma_f32_32x32x16_bf16(a, a, result, 0, 0, 0); + result = __builtin_amdgcn_mfma_f32_16x16x32_bf16(a, a, result, 0, 0, 0); } if (result[0] != 2*result[0]) From fceff8096faaa75d55a5e219f0348ac3020a27ca Mon Sep 17 00:00:00 2001 From: Vignesh Edithal Date: Fri, 16 Oct 2026 16:23:36 +0000 Subject: [PATCH 42/45] [rocprofiler-compute] Remove unused roofline benchmark constants Co-Authored-By: Claude Opus 5.5 (1M context) --- .../src/roofline/benchmark/benchmark_base.py | 2 -- 1 file changed, 2 deletions(-) diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py index 4e1efb05ca16..8d4d34ad9c06 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py @@ -38,8 +38,6 @@ # ============================================================================= DEFAULT_WORKGROUP_SIZE = 256 -DEFAULT_WORKGROUPS = 8192 -DEFAULT_THREADS = DEFAULT_WORKGROUP_SIZE * DEFAULT_WORKGROUPS DEFAULT_NUM_EXPERIMENTS = 10 DEFAULT_NUM_ITERS = 10 From 7171587bbb0ec0a123b3e15e98036d676602cde7 Mon Sep 17 00:00:00 2001 From: Vignesh Edithal Date: Fri, 16 Oct 2026 16:26:21 +0000 Subject: [PATCH 43/45] [rocprofiler-compute] Note higher gfx950 FP16/BF16 roofline peaks in CHANGELOG Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/rocprofiler-compute/CHANGELOG.md | 2 ++ 1 file changed, 2 insertions(+) diff --git a/projects/rocprofiler-compute/CHANGELOG.md b/projects/rocprofiler-compute/CHANGELOG.md index 460edb3a2dfb..99cb33fe41f9 100644 --- a/projects/rocprofiler-compute/CHANGELOG.md +++ b/projects/rocprofiler-compute/CHANGELOG.md @@ -21,6 +21,8 @@ Full documentation for ROCm Compute Profiler is available at [https://rocm.docs. * The roofline benchmark now reports the peak measured sample across runs instead of the mean. +* The roofline benchmark now reports higher FP16 and BF16 MFMA peaks on gfx950 (MI350 series), closer to the hardware's theoretical peak. + * Renamed the Memory Chart left-hand panel from "Kernel" to "Compute Units" on CDNA (gfx9) and to "WGPs" on gfx115x and gfx1250. The panel now shows resource allocation stats: Wave Occupancy as a percent of the maximum wavefronts per CU (gfx9) or per WGP (gfx115x, gfx1250), vGPRs, sGPRs, Scratch in KB per wave, LDS Allocation, and Workgroups per active CU (gfx9) or per WGP (gfx115x, gfx1250). * All Memory Chart bandwidth values now use uniform fixed-point GB/s formatting (3 decimal places) for easy cross-level comparison. From 6481efe38e623fff173df50ca04aa84b84470e71 Mon Sep 17 00:00:00 2001 From: Vignesh Edithal Date: Fri, 16 Oct 2026 16:38:37 +0000 Subject: [PATCH 44/45] [rocprofiler-compute] Tidy roofline review follow-ups - Indent the Limited by docs line so the AI axis bullet list renders - Use max(samples) directly instead of calc_stats - Move roofline_main test imports to module level - Replace the nested ternary in exportViewSubtitle with peakLabel Co-Authored-By: Claude Opus 5.5 (1M context) --- .../docs/how-to/analyze/cli.rst | 2 +- .../src/roofline/assets/roofline_plot.js | 19 +++++++++++-------- .../src/roofline/benchmark/benchmark_base.py | 14 +++++--------- .../tests/unit/roofline/test_roofline_main.py | 14 +++----------- 4 files changed, 20 insertions(+), 29 deletions(-) diff --git a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst index 6ef8fe06170e..4755205a4dcb 100644 --- a/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst +++ b/projects/rocprofiler-compute/docs/how-to/analyze/cli.rst @@ -625,7 +625,7 @@ Example with multiple ``--mem-level`` and ``--roofline-data-type`` options: Interactive Roofline HTML: * Use the *AI axis* selector to choose one memory level per kernel, *All peaks* to plot every level at once, or *Limited by* to show each kernel's limiting memory-level point (falling back to its leftmost point when no memory-level limiter is available). -Isolating one kernel shows it across all available memory levels. + Isolating one kernel shows it across all available memory levels. * Use the *Precision* selector to choose one or more arithmetic precision's peak roofline to display on the plot. * Use the *Kernels* and *Bandwidth rooflines* panels to isolate, multi-select, or reset plotted items. The *Runtime shown* slider filters to the heaviest kernels that reach the selected GPU resident-time cutoff. * Hover over kernel dots and rooflines to see arithmetic intensity, throughput, roofline percentage, limiter, runtime, bandwidth, and compute-peak details. diff --git a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js index 0cc17ddbd83d..f7e10cc91bf8 100644 --- a/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js +++ b/projects/rocprofiler-compute/src/roofline/assets/roofline_plot.js @@ -1189,15 +1189,18 @@ return entries; } + function peakLabel(peak) { + if (peak === ALL_PEAKS_VALUE) { + return ALL_PEAKS_LABEL; + } + if (peak === LIMITING_PEAK_VALUE) { + return LIMITING_PEAK_LABEL; + } + return peak; + } + function exportViewSubtitle() { - var peak = effectivePeak(); - var peakLabel = - peak === ALL_PEAKS_VALUE - ? ALL_PEAKS_LABEL - : peak === LIMITING_PEAK_VALUE - ? LIMITING_PEAK_LABEL - : peak; - var parts = ["AI axis: " + peakLabel]; + var parts = ["AI axis: " + peakLabel(effectivePeak())]; if (hasRuntimeData && isFinite(state.runtimeThreshold)) { parts.push("runtime shown: " + state.runtimeThreshold.toFixed(3) + "%"); } diff --git a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py index 8d4d34ad9c06..6dc76b998064 100644 --- a/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py +++ b/projects/rocprofiler-compute/src/roofline/benchmark/benchmark_base.py @@ -160,10 +160,6 @@ def show_progress(self, pct: float) -> None: print(f"\r{int(pct * 100):3d}% {bar}", end="", flush=True) - def calc_stats(self, samples: list) -> float: - """Returns the peak (max) value from the samples.""" - return max(samples) - class Program: """Helper class for loading and compiling kernels.""" @@ -444,7 +440,7 @@ def hbm_bw_benchmark(self, device: int) -> float: [d_src, c_int64(num_steps)], ) - peak = self.calc_stats(samples) + peak = max(samples) event_ms = total_bytes / peak / 1e6 @@ -499,7 +495,7 @@ def cache_bw_bench(self, device: int, type: str, iters: int) -> float: [mem_block, dummy, iters], ) - peak = self.calc_stats(samples) + peak = max(samples) event_ms = total_bytes / peak / 1e6 @@ -562,7 +558,7 @@ def lds_bw_benchmark(self, device: int) -> float: [iters, dummy], ) - peak = self.calc_stats(samples) + peak = max(samples) event_ms = total_bytes / peak / 1e6 @@ -625,7 +621,7 @@ def flops_bench(self, device: int, type: str, unit: str, rate: int) -> float: [memblock, iterations], ) - peak = self.calc_stats(samples) + peak = max(samples) event_ms = total_flops / peak / 1e6 @@ -690,7 +686,7 @@ def matrix_bench(self, device: int, type: str, unit: str, rate: int) -> float: [iters, dummy], ) - peak = self.calc_stats(samples) + peak = max(samples) event_ms = total_flops / peak / 1e6 diff --git a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py index c5f4ebaaf0a6..db6a28edd473 100644 --- a/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py +++ b/projects/rocprofiler-compute/tests/unit/roofline/test_roofline_main.py @@ -11,7 +11,7 @@ import math import re from pathlib import Path -from typing import TYPE_CHECKING, Dict, List, Tuple +from typing import Dict, List, Tuple import plotly.graph_objects as go import pytest @@ -20,9 +20,7 @@ from roofline.roofline_frame import FRAME_X_MIN, canonical_frame from roofline.roofline_hover import truncate_kernel_name, wrap_hover_name from roofline.roofline_html import RooflineViewModel, build_interactive_document - -if TYPE_CHECKING: - from roofline.roofline_main import Roofline +from roofline.roofline_main import Roofline, _default_precisions _ASSETS = Path(roofline_html.__file__).parent / "assets" @@ -36,11 +34,9 @@ class MockMspec: gpu_arch = "gfx90a" -def make_roofline(datatypes: List[str], **run_parameters: object) -> "Roofline": +def make_roofline(datatypes: List[str], **run_parameters: object) -> Roofline: """A Roofline for the unit tests. It never reads its ``args`` on the cli_generate_plot / generate_plot paths, so a bare Namespace suffices.""" - from roofline.roofline_main import Roofline - parameters: Dict[str, object] = { "workload_dir": "", "device_id": 0, @@ -717,8 +713,6 @@ def test_dash_figures_keep_every_ceiling(benchmarked_roofline) -> None: def test_default_precisions_fp32_fp16_fp8_present__orders_fp32_first() -> None: - from roofline.roofline_main import _default_precisions - assert _default_precisions(["FP16", "FP32", "FP8", "FP64"]) == [ "FP32", "FP16", @@ -728,8 +722,6 @@ def test_default_precisions_fp32_fp16_fp8_present__orders_fp32_first() -> None: def test_default_precisions_no_preferred__falls_back_to_first_datatype() -> None: - from roofline.roofline_main import _default_precisions - assert _default_precisions(["FP64", "BF16"]) == ["FP64"] From 5630f67367cc01932bf950ce1ddc3d2c46be5ef3 Mon Sep 17 00:00:00 2001 From: Vignesh Edithal Date: Fri, 16 Oct 2026 16:50:55 +0000 Subject: [PATCH 45/45] [rocprofiler-compute] Use every drawn datatype as a roofline limiter candidate Co-Authored-By: Claude Opus 5.5 (1M context) --- projects/rocprofiler-compute/src/roofline/roofline_main.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/rocprofiler-compute/src/roofline/roofline_main.py b/projects/rocprofiler-compute/src/roofline/roofline_main.py index fa3b753740fe..69bb12fa3423 100644 --- a/projects/rocprofiler-compute/src/roofline/roofline_main.py +++ b/projects/rocprofiler-compute/src/roofline/roofline_main.py @@ -341,7 +341,7 @@ def _figure_compute_peaks(self, ops_flops: str) -> list[tuple[str, float]]: def _collect_compute_peaks(self, ops_flops: str) -> list[tuple[str, float]]: """Read every stacked datatype's ceilings and label each peak with it.""" peaks: list[tuple[str, float]] = [] - for dt in self.__run_parameters.get("roofline_data_type", []): + for dt in SUPPORTED_DATATYPES.get(self.__mspec.gpu_arch, {}): dt = str(dt) if _figure_class(dt) != ops_flops or not self._datatype_supported(dt): continue