Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 17 additions & 2 deletions src/accordo/snapshot.py
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
"""Snapshot: Represents captured kernel argument data from a binary execution."""

from dataclasses import dataclass
from typing import List
from typing import List, Optional

import numpy as np

Expand All @@ -18,13 +18,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")
"""
Expand All @@ -33,6 +37,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."""
Expand All @@ -54,6 +60,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}")

Expand Down
35 changes: 35 additions & 0 deletions src/accordo/src/accordo.hip
Original file line number Diff line number Diff line change
Expand Up @@ -45,6 +45,7 @@ SOFTWARE.
#include <memory>
#include <optional>
#include <sstream>
#include <fstream>

#include <hip/hip_runtime.h>
#include "KernelArguments.hpp"
Expand Down Expand Up @@ -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<const hsa_kernel_dispatch_packet_t*>(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<uint32_t>(disp->grid_size_x) << ","
<< "\"y\":" << static_cast<uint32_t>(disp->grid_size_y) << ","
<< "\"z\":" << static_cast<uint32_t>(disp->grid_size_z) << "},";
ofs << "\"block\":{"
<< "\"x\":" << static_cast<uint32_t>(disp->workgroup_size_x) << ","
<< "\"y\":" << static_cast<uint32_t>(disp->workgroup_size_y) << ","
<< "\"z\":" << static_cast<uint32_t>(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);

Expand Down
35 changes: 33 additions & 2 deletions src/accordo/validator.py
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,7 @@

"""AccordoValidator: Main validation class for Accordo."""

import json
import logging
import os
import signal
Expand Down Expand Up @@ -199,17 +200,39 @@ 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)
Expand Down Expand Up @@ -298,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
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.

Expand All @@ -307,6 +335,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
Expand All @@ -324,6 +353,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()
Expand Down
Loading