From 4430f1e2d08a84129858083e70978b7df0f26062 Mon Sep 17 00:00:00 2001 From: coleramos425 Date: Fri, 21 Nov 2025 16:46:34 -0600 Subject: [PATCH 1/2] Add grid and block information for kernel arguments Signed-off-by: coleramos425 --- src/accordo/snapshot.py | 18 +++++++++++++++++- src/accordo/src/accordo.hip | 35 +++++++++++++++++++++++++++++++++++ src/accordo/validator.py | 26 ++++++++++++++++++++++++-- 3 files changed, 76 insertions(+), 3 deletions(-) diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py index 025f5020..44f72266 100644 --- a/src/accordo/snapshot.py +++ b/src/accordo/snapshot.py @@ -5,6 +5,7 @@ from dataclasses import dataclass from typing import List +from typing import Optional import numpy as np @@ -18,13 +19,17 @@ class Snapshot: execution_time_ms: Time taken to execute and capture the snapshot (milliseconds) binary: The binary command that was executed working_directory: The directory where the binary was executed + grid_size: Optional grid dimensions dict with x,y,z (if available) + block_size: Optional workgroup dimensions dict with x,y,z (if available) Example: >>> snapshot = Snapshot( ... arrays=[np.array([1, 2, 3]), np.array([4, 5, 6])], ... execution_time_ms=12.5, ... binary=["./my_app"], - ... working_directory="/path/to/project" + ... working_directory="/path/to/project", + ... grid_size={"x": 1, "y": 1, "z": 1}, + ... block_size={"x": 256, "y": 1, "z": 1}, ... ) >>> print(f"Captured {len(snapshot.arrays)} arrays in {snapshot.execution_time_ms}ms") """ @@ -33,6 +38,8 @@ class Snapshot: execution_time_ms: float binary: List[str] working_directory: str + grid_size: Optional[dict] = None + block_size: Optional[dict] = None def __repr__(self) -> str: """Pretty representation of snapshot.""" @@ -54,6 +61,15 @@ def summary(self) -> str: f" Number of Arrays: {len(self.arrays)}", ] + if self.grid_size is not None: + lines.append( + f" Grid Size: x={self.grid_size.get('x')}, y={self.grid_size.get('y')}, z={self.grid_size.get('z')}" + ) + if self.block_size is not None: + lines.append( + f" Block Size: x={self.block_size.get('x')}, y={self.block_size.get('y')}, z={self.block_size.get('z')}" + ) + for i, arr in enumerate(self.arrays): lines.append(f" Array {i}: shape={arr.shape}, dtype={arr.dtype}") diff --git a/src/accordo/src/accordo.hip b/src/accordo/src/accordo.hip index 3aad7220..4f0d916f 100644 --- a/src/accordo/src/accordo.hip +++ b/src/accordo/src/accordo.hip @@ -45,6 +45,7 @@ SOFTWARE. #include #include #include +#include #include #include "KernelArguments.hpp" @@ -503,6 +504,40 @@ void accordo::write_packets(hsa_queue_t* queue, LOG_DETAIL("Executing packet: {}", packet_to_text(packet)); auto instance = get_instance(); + // Best-effort: write dispatch metadata (grid/block dims) once if env is set and kernel matches + static bool wrote_metadata = false; + if (!wrote_metadata) { + uint32_t type = get_header_type(packet); + if (type == HSA_PACKET_TYPE_KERNEL_DISPATCH) { + const hsa_kernel_dispatch_packet_t* disp = + reinterpret_cast(packet); + const auto kernel_name = get_kernel_name(disp->kernel_object); + static const char* kernel_to_trace = std::getenv("KERNEL_TO_TRACE"); + static const char* metadata_file = std::getenv("ACCORDO_METADATA_FILE"); + if (metadata_file && kernel_to_trace && kernel_name.contains(kernel_to_trace)) { + std::ofstream ofs(metadata_file, std::ios::out | std::ios::trunc); + if (ofs) { + ofs << "{"; + ofs << "\"kernel_name\":\"" << kernel_name << "\","; + ofs << "\"grid\":{" + << "\"x\":" << static_cast(disp->grid_size_x) << "," + << "\"y\":" << static_cast(disp->grid_size_y) << "," + << "\"z\":" << static_cast(disp->grid_size_z) << "},"; + ofs << "\"block\":{" + << "\"x\":" << static_cast(disp->workgroup_size_x) << "," + << "\"y\":" << static_cast(disp->workgroup_size_y) << "," + << "\"z\":" << static_cast(disp->workgroup_size_z) << "}"; + ofs << "}"; + ofs.flush(); + wrote_metadata = true; + LOG_DETAIL("Wrote dispatch metadata to {}", metadata_file); + } else { + LOG_ERROR("Failed to open metadata file {}", metadata_file); + } + } + } + } + hsa_signal_t new_signal; auto status = hsa_core_call(instance, hsa_signal_create, 1, 0, nullptr, &new_signal); diff --git a/src/accordo/validator.py b/src/accordo/validator.py index bb77f085..9a3f2ab6 100644 --- a/src/accordo/validator.py +++ b/src/accordo/validator.py @@ -10,6 +10,7 @@ import time from pathlib import Path from typing import Optional +import json import numpy as np @@ -199,17 +200,35 @@ def capture_snapshot( try: start_time = time.time() + # Prepare a metadata file path for exporter to write dispatch dims + metadata_file = f"/tmp/accordo_metadata_{int(start_time * 1000)}.json" result_arrays = self._run_instrumented_app( - binary, working_directory, label="snapshot", baseline_time_ms=None + binary, working_directory, label="snapshot", baseline_time_ms=None, extra_env={"ACCORDO_METADATA_FILE": metadata_file} ) signal.alarm(0) # Cancel alarm on success execution_time_ms = (time.time() - start_time) * 1000 + # Try to read grid/block metadata if exporter produced it + grid = None + block = None + try: + if os.path.exists(metadata_file): + with open(metadata_file, "r") as f: + meta = json.load(f) + # Basic validation + if isinstance(meta, dict): + grid = meta.get("grid") + block = meta.get("block") + except Exception: + # TODO: WIP; current attempt doesn't catch parsing errors + pass return Snapshot( arrays=result_arrays, execution_time_ms=execution_time_ms, binary=binary, working_directory=working_directory, + grid_size=grid, + block_size=block, ) except _TimeoutException: signal.alarm(0) @@ -298,7 +317,7 @@ def validate( return self.compare_snapshots(reference_snapshot, optimized_snapshot) def _run_instrumented_app( - self, binary_cmd: list[str], working_directory: str, label: str, baseline_time_ms: Optional[float] = None + self, binary_cmd: list[str], working_directory: str, label: str, baseline_time_ms: Optional[float] = None, extra_env: Optional[dict] = None ) -> list[np.ndarray]: """Run an instrumented application and collect kernel argument data. @@ -307,6 +326,7 @@ def _run_instrumented_app( working_directory: Directory to run the binary from label: Label for this run ("reference" or "optimized") baseline_time_ms: Baseline time for dynamic timeout + extra_env: Optional environment variables to inject for this run Returns: List of numpy arrays with kernel argument data @@ -324,6 +344,8 @@ def _run_instrumented_app( env = os.environ.copy() env["HSA_TOOLS_LIB"] = str(self._lib_path) env["KERNEL_TO_TRACE"] = self.config.kernel_name + if extra_env: + env.update(extra_env) # Set log level debug_level = logging.getLogger().getEffectiveLevel() From 7be45968d2be9b4d78d0fb4071ccb3b3dda5f47e Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" Date: Fri, 21 Nov 2025 22:46:53 +0000 Subject: [PATCH 2/2] Apply Ruff auto-fixes --- src/accordo/snapshot.py | 3 +-- src/accordo/validator.py | 15 ++++++++++++--- 2 files changed, 13 insertions(+), 5 deletions(-) diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py index 44f72266..dc716fbd 100644 --- a/src/accordo/snapshot.py +++ b/src/accordo/snapshot.py @@ -4,8 +4,7 @@ """Snapshot: Represents captured kernel argument data from a binary execution.""" from dataclasses import dataclass -from typing import List -from typing import Optional +from typing import List, Optional import numpy as np diff --git a/src/accordo/validator.py b/src/accordo/validator.py index 9a3f2ab6..06a08f5c 100644 --- a/src/accordo/validator.py +++ b/src/accordo/validator.py @@ -3,6 +3,7 @@ """AccordoValidator: Main validation class for Accordo.""" +import json import logging import os import signal @@ -10,7 +11,6 @@ import time from pathlib import Path from typing import Optional -import json import numpy as np @@ -203,7 +203,11 @@ def capture_snapshot( # Prepare a metadata file path for exporter to write dispatch dims metadata_file = f"/tmp/accordo_metadata_{int(start_time * 1000)}.json" result_arrays = self._run_instrumented_app( - binary, working_directory, label="snapshot", baseline_time_ms=None, extra_env={"ACCORDO_METADATA_FILE": metadata_file} + binary, + working_directory, + label="snapshot", + baseline_time_ms=None, + extra_env={"ACCORDO_METADATA_FILE": metadata_file}, ) signal.alarm(0) # Cancel alarm on success execution_time_ms = (time.time() - start_time) * 1000 @@ -317,7 +321,12 @@ def validate( return self.compare_snapshots(reference_snapshot, optimized_snapshot) def _run_instrumented_app( - self, binary_cmd: list[str], working_directory: str, label: str, baseline_time_ms: Optional[float] = None, extra_env: Optional[dict] = None + self, + binary_cmd: list[str], + working_directory: str, + label: str, + baseline_time_ms: Optional[float] = None, + extra_env: Optional[dict] = None, ) -> list[np.ndarray]: """Run an instrumented application and collect kernel argument data.