From 0937a78de6bc737c0a5353e897b985693a5a0a6c Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 00:49:19 -0600 Subject: [PATCH 01/15] Add Accordo API refactoring TODO --- accordo.todo | 334 +++++++++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 334 insertions(+) create mode 100644 accordo.todo diff --git a/accordo.todo b/accordo.todo new file mode 100644 index 00000000..edbca5e4 --- /dev/null +++ b/accordo.todo @@ -0,0 +1,334 @@ +# Accordo API Refactoring TODO + +## Current Problems +- [ ] No clear abstraction layers - CMake builds, pipe management, IPC handling all exposed +- [ ] Scattered concerns - code_gen, communicate, utils in separate modules +- [ ] Hard to use - Formula authors must manage pipes, env vars, processes manually +- [ ] Poor testability - Everything coupled to formula_base +- [ ] No clean imports - `from accordo.python.code_gen import generate_header` is ugly +- [ ] Strings for kernel args - no semantic info, bad for LLM generation + +## Proposed Clean API Design + +### High-Level User API (3 Main Classes) + +```python +from accordo import AccordoValidator, ValidationConfig, ValidationResult + +# Clean, simple usage +config = ValidationConfig( + kernel_name="my_kernel", + kernel_args=[ + KernelArg(name="result", type="double*", direction="out"), + KernelArg(name="input", type="const double*", direction="in"), + KernelArg(name="count", type="unsigned long", direction="in") + ], + additional_includes=['"my_types.h"'], # For custom types + tolerance=1e-6, + timeout_multiplier=2.0 +) + +validator = AccordoValidator(config) +result = validator.validate(reference_app, optimized_app, baseline_time_ms=10.5) + +if result.is_valid: + print(f"✓ Validation passed! {result.num_arrays_validated} arrays matched") +else: + print(f"✗ Failed: {result.error_message}") + for m in result.mismatches: + print(f" Arg '{m.arg_name}' ({m.arg_type}): max_diff={m.max_difference}") +``` + +### Package Structure + +``` +src/accordo/ +├── __init__.py # Public API: AccordoValidator, ValidationConfig, ValidationResult +├── validator.py # AccordoValidator class +├── config.py # ValidationConfig, KernelArg dataclasses +├── result.py # ValidationResult, ArrayMismatch dataclasses +├── exceptions.py # AccordoError, AccordoBuildError, AccordoTimeoutError, etc. +├── _internal/ # Private implementation (not imported by users) +│ ├── __init__.py +│ ├── builder.py # AccordoBuilder - manages CMake builds +│ ├── runtime.py # AccordoRuntime - manages process + IPC +│ ├── codegen.py # generate_kernel_header() - creates KernelArguments.hpp +│ ├── ipc/ +│ │ ├── __init__.py +│ │ ├── communication.py # IPC handle reading/writing +│ │ ├── memory.py # Device-to-host memory operations +│ │ └── types.py # Type mappings (float*, int*, etc) +│ └── hip.py # HIP interop functions +└── cpp/ # C++ library (unchanged) + ├── accordo.hpp + ├── accordo.hip + └── CMakeLists.txt +``` + +## Key Design Decisions + +### 1. Structured KernelArg (Not Plain Strings) +**Rationale:** LLMs generate JSON better than parsing strings. Named args give better error messages. + +```python +# ✅ Structured (LLM-friendly, semantic, extensible) +kernel_args=[ + KernelArg(name="result", type="double*", direction="out"), + KernelArg(name="input", type="const double*", direction="in") +] + +# ❌ Old way (hard for LLM, unclear semantics) +kernel_args=["double*", "const double*"] +``` + +**Benefits:** +- LLMs excel at generating structured JSON +- Better error messages: "Arg 'result' failed" vs "Arg 0 failed" +- Explicit input/output direction +- Easy to extend with metadata (size, validate flag, etc) +- Backward compat: still accept strings, auto-convert to KernelArg + +### 2. No BuildConfig Class +**Rationale:** Build is deterministic, fast, and users don't care about build details. + +- Accordo auto-builds on first use +- CMake caching makes rebuilds instant +- Optional kwargs for power users: `force_rebuild=True`, `parallel_jobs=16` + +### 3. Flat Arguments Only (For Now) +**Rationale:** 95% of GPU kernels use flat signatures. Nested struct support can be added later. + +```python +# ✅ Supported (flat args) +kernel_args=[ + KernelArg("output", "float*"), + KernelArg("input", "const float*"), + KernelArg("size", "int") +] + +# ❌ Not supported yet (nested) +kernel_args=[ + KernelArg("matrix", "Matrix*") # Matrix has nested pointers +] + +# Future: Add visitor/accessor pattern for nested types +``` + +### 4. additional_includes for Custom Types +**Rationale:** Accordo generates C++ header, needs includes to define custom types. + +```python +config = ValidationConfig( + kernel_name="my_kernel", + kernel_args=[ + KernelArg("matrix", "MyMatrix*"), + KernelArg("data", "__hip_bfloat16*") + ], + additional_includes=[ + '"my_app/matrix.h"', # Custom types + '' # HIP types + ] +) +``` + +Generated header: +```cpp +#pragma once +#include +#include "my_app/matrix.h" // User's custom types +#include // HIP types + +struct KernelArguments { + MyMatrix* arg0; + __hip_bfloat16* arg1; + + auto as_tuple() const { + return std::tie(arg0, arg1); + } +}; +``` + +## Core Classes Detail + +### ValidationConfig +```python +@dataclass +class KernelArg: + name: str # "result", "input", "count" + type: str # "double*", "const float*", "int" + direction: str = "auto" # "in", "out", "inout", "auto" (infer from const) + +@dataclass +class ValidationConfig: + kernel_name: str + kernel_args: list[Union[KernelArg, str, dict]] # Accepts multiple formats + additional_includes: list[str] = field(default_factory=list) + tolerance: float = 1e-6 + timeout_multiplier: float = 2.0 + log_level: str = "WARNING" +``` + +### ValidationResult +```python +@dataclass +class ArrayMismatch: + arg_index: int + arg_name: str + arg_type: str + max_difference: float + mean_difference: float + reference_sample: np.ndarray + optimized_sample: np.ndarray + +@dataclass +class ValidationResult: + is_valid: bool + error_message: Optional[str] + mismatches: list[ArrayMismatch] + matched_arrays: dict[str, dict] + execution_time_ms: dict[str, float] # {"reference": 10.5, "optimized": 8.3} + timeout_used: Optional[float] + + @property + def num_arrays_validated(self) -> int: + return len(self.matched_arrays) + len(self.mismatches) +``` + +### AccordoValidator +```python +class AccordoValidator: + def __init__( + self, + config: ValidationConfig, + accordo_path: Optional[Path] = None, # Auto-detected + force_rebuild: bool = False, + parallel_jobs: int = 16 + ): + """Initialize validator. Lazy builds on first validate() call.""" + + def validate( + self, + reference_app: Application, + optimized_app: Application, + baseline_time_ms: Optional[float] = None + ) -> ValidationResult: + """ + Validate optimized vs reference. + + Args: + baseline_time_ms: For dynamic timeout (timeout = baseline * multiplier) + + Returns: + ValidationResult with validation status and details + + Raises: + AccordoTimeoutError, AccordoProcessError, AccordoBuildError + """ +``` + +### Exceptions +```python +class AccordoError(Exception): pass +class AccordoBuildError(AccordoError): pass +class AccordoTimeoutError(AccordoError): + def __init__(self, message, timeout_seconds): + self.timeout_seconds = timeout_seconds +class AccordoProcessError(AccordoError): + def __init__(self, message, exit_code=None): + self.exit_code = exit_code +``` + +## Integration with IntelliPerf + +### Simplified formula_base.py +```python +from accordo import AccordoValidator, ValidationConfig, KernelArg + +def correctness_validation_pass(self, kernel, kernel_args, tolerance=1e-6): + # Create config + config = ValidationConfig( + kernel_name=kernel, + kernel_args=kernel_args, # List of KernelArg or strings + tolerance=tolerance + ) + + # Validate + validator = AccordoValidator(config) + result = validator.validate( + reference_app=self._reference_app, + optimized_app=self._application, + baseline_time_ms=getattr(self, "baseline_time_ms", None) + ) + + return Result( + success=result.is_valid, + error_report=result.error_message if not result.is_valid else "" + ) +``` + +## Implementation Tasks + +- [ ] Create new package structure +- [ ] Implement config.py (ValidationConfig, KernelArg) +- [ ] Implement result.py (ValidationResult, ArrayMismatch) +- [ ] Implement exceptions.py (all custom exceptions) +- [ ] Implement validator.py (AccordoValidator main class) +- [ ] Implement _internal/builder.py (CMake build management) +- [ ] Implement _internal/runtime.py (process + IPC management) +- [ ] Implement _internal/codegen.py (header generation with additional_includes) +- [ ] Implement _internal/ipc/ (communication, memory, types) +- [ ] Update __init__.py (clean public API exports) +- [ ] Update formula_base.py to use new API +- [ ] Write tests for each component +- [ ] Update documentation +- [ ] Migration guide for existing code + +## Future Enhancements (Post-Flat-Args) + +### Nested Struct Support +```python +# Option 1: User provides accessor +kernel_args=[ + KernelArg( + name="matrix", + type="Matrix*", + nested_pointers=["Matrix.data"], # Specify nested pointers to validate + element_type="float*" + ) +] + +# Option 2: Type definitions for reflection +type_definitions={ + "Matrix": { + "fields": [ + {"name": "data", "type": "float*", "validate": True}, + {"name": "rows", "type": "int", "validate": False} + ] + } +} +``` + +### Size Expressions +```python +kernel_args=[ + KernelArg(name="output", type="float*", size_expr="N"), + KernelArg(name="input", type="float*", size_expr="N*M") +] +``` + +### Selective Validation +```python +kernel_args=[ + KernelArg(name="output", type="float*", validate=True), + KernelArg(name="config", type="Config*", validate=False) # Skip complex structs +] +``` + +## Notes +- Keep API minimal and clean - no BuildConfig +- Optimize for LLM generation - structured KernelArg over strings +- Start with flat args (95% use case), add nested later +- Auto-build, auto-detect, minimize user configuration +- Clear error messages with named arguments +- Type hints everywhere for IDE support + From aa25362c484167a129f44f5889c50df80cb44b96 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 00:54:28 -0600 Subject: [PATCH 02/15] WIP: Accordo API refactoring - Phase 1 Implemented core public API classes and internal modules: Public API: - config.py: ValidationConfig and KernelArg classes * Structured kernel arguments with name, type, direction * Backward compatible with strings and dicts * Support for additional_includes for custom types - result.py: ValidationResult and ArrayMismatch classes * Detailed validation results with metrics * Human-readable summary() method - exceptions.py: Custom exception hierarchy * AccordoError, AccordoBuildError, AccordoTimeoutError * AccordoProcessError, AccordoValidationError Internal modules (_internal/): - codegen.py: Header generation with additional_includes support - hip.py: HIP interop (open_ipc_handle, memcpy_d2h) - ipc/communication.py: IPC communication with dynamic timeout Updated __init__.py to export clean public API. Added README_REFACTORING.md documenting progress. TODO: - Implement AccordoValidator class - Builder and Runtime modules - Integration with formula_base.py - Tests See accordo.todo for complete design. --- src/accordo/README_REFACTORING.md | 168 +++++++++++++++++++++ src/accordo/__init__.py | 72 ++++++--- src/accordo/_internal/__init__.py | 2 + src/accordo/_internal/codegen.py | 52 +++++++ src/accordo/_internal/hip.py | 107 +++++++++++++ src/accordo/_internal/ipc/__init__.py | 2 + src/accordo/_internal/ipc/communication.py | 166 ++++++++++++++++++++ src/accordo/config.py | 120 +++++++++++++++ src/accordo/exceptions.py | 36 +++++ src/accordo/result.py | 101 +++++++++++++ 10 files changed, 802 insertions(+), 24 deletions(-) create mode 100644 src/accordo/README_REFACTORING.md create mode 100644 src/accordo/_internal/__init__.py create mode 100644 src/accordo/_internal/codegen.py create mode 100644 src/accordo/_internal/hip.py create mode 100644 src/accordo/_internal/ipc/__init__.py create mode 100644 src/accordo/_internal/ipc/communication.py create mode 100644 src/accordo/config.py create mode 100644 src/accordo/exceptions.py create mode 100644 src/accordo/result.py diff --git a/src/accordo/README_REFACTORING.md b/src/accordo/README_REFACTORING.md new file mode 100644 index 00000000..a788e38e --- /dev/null +++ b/src/accordo/README_REFACTORING.md @@ -0,0 +1,168 @@ +# Accordo API Refactoring Progress + +## ✅ Completed + +### 1. Package Structure +- Created `_internal/` directory for private implementation +- Created `_internal/ipc/` subdirectory for IPC modules +- Added proper `__init__.py` files + +### 2. Public API Classes + +#### `config.py` - Configuration +- ✅ `KernelArg`: Structured kernel argument with name, type, and direction + - Supports auto-inference of direction from type + - Backward compatible with plain strings and dicts +- ✅ `ValidationConfig`: Configuration for validation + - Supports `kernel_args` as KernelArg, str, or dict + - Includes `additional_includes` for custom types + - Configurable tolerance and timeout + +#### `result.py` - Results +- ✅ `ArrayMismatch`: Detailed mismatch information +- ✅ `ValidationResult`: Validation result with metrics + - Properties: `num_arrays_validated`, `success_rate` + - Method: `summary()` for human-readable output + +#### `exceptions.py` - Exceptions +- ✅ `AccordoError`: Base exception +- ✅ `AccordoBuildError`: Build failures +- ✅ `AccordoTimeoutError`: Timeout errors (with timeout_seconds attribute) +- ✅ `AccordoProcessError`: Process crashes (with exit_code attribute) +- ✅ `AccordoValidationError`: Validation failures + +### 3. Internal Modules + +#### `_internal/codegen.py` +- ✅ `generate_kernel_header()`: Creates KernelArguments.hpp + - Supports `additional_includes` parameter + - Generates flat argument structures + +#### `_internal/hip.py` +- ✅ `hip_try()`: HIP error checking +- ✅ `open_ipc_handle()`: Open IPC memory handles +- ✅ `memcpy_d2h()`: Device-to-host memory copy + +#### `_internal/ipc/communication.py` +- ✅ `read_ipc_handles()`: Read IPC handles from file +- ✅ `send_response()`: Send completion signal via pipe +- ✅ `get_kern_arg_data()`: Get kernel argument data via IPC + - Supports dynamic timeout based on baseline + - Type mapping for common GPU types (float*, double*, __half*, etc.) + +### 4. Public API Export +- ✅ Updated `__init__.py` to export public API classes +- ✅ Clean imports: `from accordo import ValidationConfig, KernelArg, ValidationResult` + +## 🚧 TODO + +### 1. AccordoValidator Implementation +The main validator class needs to be implemented. This will: +- Wrap the existing functionality from `formula_base.py` +- Manage Accordo C++ library building (via CMake) +- Handle process management and IPC communication +- Provide the `validate(reference_app, optimized_app)` method + +**Key Design Decisions:** +- Lazy build on first use +- Auto-detect Accordo path +- Dynamic timeout calculation +- Clean error handling with custom exceptions + +### 2. Builder Module (`_internal/builder.py`) +Manage CMake builds of Accordo C++ library: +- Check if build needed +- Run CMake with proper flags +- Handle build errors +- Cache build results + +### 3. Runtime Module (`_internal/runtime.py`) +Manage process execution and IPC: +- Launch instrumented processes +- Monitor process health +- Handle timeouts +- Clean up resources + +### 4. Integration with IntelliPerf +Update `formula_base.py` to use the new API: +```python +from accordo import AccordoValidator, ValidationConfig, KernelArg + +config = ValidationConfig( + kernel_name=kernel, + kernel_args=[KernelArg(name=f"arg{i}", type=t) for i, t in enumerate(kernel_args)], + tolerance=tolerance +) + +validator = AccordoValidator(config) +result = validator.validate( + reference_app=self._reference_app, + optimized_app=self._application, + baseline_time_ms=getattr(self, "baseline_time_ms", None) +) + +return Result( + success=result.is_valid, + error_report=result.error_message if not result.is_valid else "" +) +``` + +### 5. Tests +- Unit tests for each module +- Integration tests with real kernels +- Test backward compatibility + +### 6. Documentation +- API reference docs +- Usage examples +- Migration guide + +## 📝 Design Notes + +### Flat Arguments Only (For Now) +Current implementation supports only flat argument lists: +```python +# ✅ Supported +kernel_args=[ + KernelArg("output", "float*"), + KernelArg("input", "const float*"), + KernelArg("size", "int") +] + +# ❌ Not yet supported (nested structs) +kernel_args=[ + KernelArg("matrix", "Matrix*") # Matrix has nested pointers +] +``` + +Future enhancement: Add visitor/accessor pattern for nested types. + +### Backward Compatibility +The API accepts multiple input formats: +- `KernelArg` objects (new, preferred) +- Plain strings (backward compatible) +- Dicts (for JSON/LLM generation) + +All formats are automatically normalized to `KernelArg` instances. + +### Type Mapping +Supported GPU types in `_internal/ipc/communication.py`: +- `double*`, `float*` +- `int*`, `std::size_t*` +- `__half*` (FP16) +- `__hip_bfloat16*` (BFloat16) + +Additional types can be added to the `type_map` dictionary. + +## 🎯 Next Steps + +1. **Implement AccordoValidator** - This is the main missing piece +2. **Test with existing formulas** - Ensure backward compatibility +3. **Update formula_base.py** - Integrate new API +4. **Add comprehensive tests** - Cover edge cases +5. **Document migration path** - Help users adopt new API + +## 📚 Reference + +See `accordo.todo` for the complete design document. + diff --git a/src/accordo/__init__.py b/src/accordo/__init__.py index 946f447d..bc0df9d1 100644 --- a/src/accordo/__init__.py +++ b/src/accordo/__init__.py @@ -1,29 +1,53 @@ -################################################################################ -# MIT License +"""Accordo: Automated side-by-side correctness validation for GPU kernels. -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. +Public API: + - ValidationConfig: Configuration for kernel validation + - KernelArg: Structured kernel argument representation + - ValidationResult: Result of validation with detailed metrics + - ArrayMismatch: Information about array validation failures + - AccordoValidator: Main validator class (TODO: implement) + - Exceptions: AccordoError, AccordoBuildError, AccordoTimeoutError, etc. -# Permission is hereby granted, free of charge, to any person obtaining a copy -# of this software and associated documentation files (the "Software"), to deal -# in the Software without restriction, including without limitation the rights -# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -# copies of the Software, and to permit persons to whom the Software is -# furnished to do so, subject to the following conditions: +Example: + >>> from accordo import ValidationConfig, KernelArg + >>> config = ValidationConfig( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... KernelArg(name="result", type="double*", direction="out"), + ... KernelArg(name="input", type="const double*", direction="in"), + ... ], + ... tolerance=1e-6 + ... ) +""" -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. +# Public API exports +from .config import KernelArg, ValidationConfig +from .exceptions import ( + AccordoBuildError, + AccordoError, + AccordoProcessError, + AccordoTimeoutError, + AccordoValidationError, +) +from .result import ArrayMismatch, ValidationResult -# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -# SOFTWARE. -################################################################################ +# Version +__version__ = "0.2.0-dev" -"""Accordo package for validation and verification.""" - -from .python import code_gen, communicate, hip, utils - -__all__ = ["communicate", "code_gen", "utils", "hip"] +# Public API +__all__ = [ + # Config + "ValidationConfig", + "KernelArg", + # Results + "ValidationResult", + "ArrayMismatch", + # Exceptions + "AccordoError", + "AccordoBuildError", + "AccordoTimeoutError", + "AccordoProcessError", + "AccordoValidationError", + # TODO: Add AccordoValidator when implemented + # "AccordoValidator", +] diff --git a/src/accordo/_internal/__init__.py b/src/accordo/_internal/__init__.py new file mode 100644 index 00000000..edfbb039 --- /dev/null +++ b/src/accordo/_internal/__init__.py @@ -0,0 +1,2 @@ +"""Internal implementation details for Accordo. Not part of public API.""" + diff --git a/src/accordo/_internal/codegen.py b/src/accordo/_internal/codegen.py new file mode 100644 index 00000000..abbae9f9 --- /dev/null +++ b/src/accordo/_internal/codegen.py @@ -0,0 +1,52 @@ +"""Code generation for Accordo C++ header files.""" + +import logging +from pathlib import Path + + +def generate_kernel_header(args: list[str], additional_includes: list[str] = None) -> str: + """Generate C++ header file for kernel arguments. + + Args: + args: List of argument type strings (e.g., ["double*", "const float*", "int"]) + additional_includes: Optional list of additional include directives + + Returns: + Path to the generated header file + """ + if additional_includes is None: + additional_includes = [] + + header_path = "/tmp/KernelArguments.hpp" + member_names = [f"arg{i}" for i in range(len(args))] + members = ";\n ".join(f"{arg} {name}" for arg, name in zip(args, member_names)) + ";" + as_tuple_members = ", ".join(member_names) + + # Build includes section + includes_section = "#include \n" + includes_section += "#include // for float16\n" + includes_section += "#include // for bfloat16\n" + + if additional_includes: + includes_section += "\n// User-provided includes\n" + for include in additional_includes: + includes_section += f"#include {include}\n" + + header_content = f"""#pragma once +{includes_section} +struct KernelArguments {{ + {members} + + auto as_tuple() const {{ + return std::tie({as_tuple_members}); + }} +}}; +""" + + with open(header_path, "w") as header_file: + header_file.write(header_content) + + logging.debug(f"Generated header file: {header_path}") + logging.debug(f"Header content: {header_content}") + return header_path + diff --git a/src/accordo/_internal/hip.py b/src/accordo/_internal/hip.py new file mode 100644 index 00000000..b05bfc0d --- /dev/null +++ b/src/accordo/_internal/hip.py @@ -0,0 +1,107 @@ +"""HIP interop functions for Accordo.""" + +import ctypes +import logging + +import numpy as np + +rt_path = "libamdhip64.so" +hip_runtime = ctypes.cdll.LoadLibrary(rt_path) + + +def hip_try(err): + """Check HIP error code and raise exception if error occurred.""" + if err != 0: + hip_runtime.hipGetErrorString.restype = ctypes.c_char_p + error_string = hip_runtime.hipGetErrorString(ctypes.c_int(err)).decode("utf-8") + raise RuntimeError(f"HIP error code {err}: {error_string}") + + +class hipIpcMemHandle_t(ctypes.Structure): + """HIP IPC memory handle structure.""" + + _fields_ = [("reserved", ctypes.c_char * 64)] + + +def open_ipc_handle(ipc_handle_data): + """Open a HIP IPC memory handle. + + Args: + ipc_handle_data: NumPy array of uint8 with 64 elements + + Returns: + Device pointer value + """ + ptr = ctypes.c_void_p() + hipIpcMemLazyEnablePeerAccess = ctypes.c_uint(1) + hip_runtime.hipIpcOpenMemHandle.argtypes = [ + ctypes.POINTER(ctypes.c_void_p), + hipIpcMemHandle_t, + ctypes.c_uint, + ] + + if isinstance(ipc_handle_data, np.ndarray): + if ipc_handle_data.dtype != np.uint8 or ipc_handle_data.size != 64: + logging.debug(f"ipc_handle_data.size: {ipc_handle_data.size}") + raise ValueError("ipc_handle_data must be a 64-element uint8 numpy array") + ipc_handle_bytes = ipc_handle_data.tobytes() + ipc_handle_data = (ctypes.c_char * 64).from_buffer_copy(ipc_handle_bytes) + else: + raise TypeError("ipc_handle_data must be a numpy.ndarray of dtype uint8 with 64 elements") + + raw_memory = ctypes.create_string_buffer(64) + ctypes.memset(raw_memory, 0x00, 64) + ipc_handle_struct = hipIpcMemHandle_t.from_buffer(raw_memory) + ipc_handle_data_bytes = bytes(ipc_handle_data) + ctypes.memmove(raw_memory, ipc_handle_data_bytes, 64) + + logging.debug("[ipc_handle_struct]:") + for i in range(0, len(ipc_handle_data_bytes), 16): + chunk = ipc_handle_data_bytes[i : i + 16] + logging.debug(" ".join(f"{b:02x}" for b in chunk)) + + hip_try( + hip_runtime.hipIpcOpenMemHandle( + ctypes.byref(ptr), + ipc_handle_struct, + hipIpcMemLazyEnablePeerAccess, + ) + ) + + return ptr.value + + +def memcpy_d2h(ptr, num_elements_to_copy, dtype): + """Copy data from device to host. + + Args: + ptr: Device pointer value + num_elements_to_copy: Number of elements to copy + dtype: C type of elements + + Returns: + NumPy array with copied data + """ + host_array = np.zeros(num_elements_to_copy, dtype=np.dtype(dtype)) + + hip_runtime.hipMemcpy.argtypes = [ + ctypes.c_void_p, + ctypes.c_void_p, + ctypes.c_size_t, + ctypes.c_int, + ] + bytes_to_copy = num_elements_to_copy * ctypes.sizeof(dtype) + logging.debug( + f"Copying {num_elements_to_copy * ctypes.sizeof(dtype)} bytes from {hex(ptr)} to {hex(host_array.ctypes.data)}" + ) + + hip_try( + hip_runtime.hipMemcpy( + ctypes.c_void_p(host_array.ctypes.data), + ctypes.c_void_p(ptr), + ctypes.c_size_t(bytes_to_copy), + 2, # hipMemcpyDeviceToHost + ) + ) + return host_array + diff --git a/src/accordo/_internal/ipc/__init__.py b/src/accordo/_internal/ipc/__init__.py new file mode 100644 index 00000000..9af22b27 --- /dev/null +++ b/src/accordo/_internal/ipc/__init__.py @@ -0,0 +1,2 @@ +"""IPC communication modules for Accordo.""" + diff --git a/src/accordo/_internal/ipc/communication.py b/src/accordo/_internal/ipc/communication.py new file mode 100644 index 00000000..200c8a2a --- /dev/null +++ b/src/accordo/_internal/ipc/communication.py @@ -0,0 +1,166 @@ +"""IPC communication for Accordo.""" + +import ctypes +import logging +import os +import stat +import time + +import ml_dtypes +import numpy as np + +from ..hip import memcpy_d2h, open_ipc_handle + + +def read_ipc_handles(args, ipc_file_name): + """Read IPC handles and sizes from the IPC file. + + Args: + args: List of argument type strings + ipc_file_name: Path to the IPC file + + Returns: + Tuple of (handles, sizes) + """ + count = sum(1 for arg in args if "*" in arg and "const" not in arg) + + handles = [] + sizes = [] + handles_set = set() + + while len(handles) < count: + if not os.path.exists(ipc_file_name): + logging.debug("Waiting for IPC file...") + time.sleep(0.1) + continue + + with open(ipc_file_name, "rb") as file: + data = file.read() + + messages = data.split(b"BEGIN\n") + for message in messages: + if b"END\n" in message: + content = message.split(b"END\n")[0] + + if len(content) == 72: + handle_data = content[:64] + size_data = content[64:72] + + handle_np = np.frombuffer(handle_data, dtype=np.uint8) + handle_tuple = tuple(handle_np) + + if handle_tuple not in handles_set: + handles.append(handle_np) + handles_set.add(handle_tuple) + + size_value = int.from_bytes(size_data, byteorder="little") + sizes.append(size_value) + + logging.debug("Final IPC Handle (hex):") + for i in range(0, len(handle_np), 16): + chunk = handle_np[i : i + 16] + logging.debug(" ".join(f"{b:02x}" for b in chunk)) + + logging.debug(f"Corresponding Pointer Size: {size_value} bytes") + + if len(handles) < count: + logging.debug(f"Waiting for {count - len(handles)} more IPC handles...") + time.sleep(0.1) + + return handles, sizes + + +def send_response(pipe_name): + """Send completion response through named pipe.""" + with open(pipe_name, "w") as fifo: + fifo.write("done\n") + + +def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30, process_pid=None, baseline_time_ms=None): + """Get kernel argument data via IPC. + + Args: + pipe_name: Path to the named pipe + args: List of argument type strings + ipc_file_name: Path to the IPC file + ipc_timeout_seconds: Timeout for IPC operations + process_pid: Process ID (for error messages) + baseline_time_ms: Baseline execution time (for dynamic timeout) + + Returns: + List of NumPy arrays with argument data + + Raises: + TimeoutError: If IPC operation times out + TypeError: If unsupported type encountered + """ + # Calculate dynamic timeout if baseline provided + if baseline_time_ms is not None: + # Use 2x baseline or minimum 3 seconds + dynamic_timeout = max(3.0, (baseline_time_ms / 1000.0) * 2.0) + ipc_timeout_seconds = dynamic_timeout + logging.debug(f"Using dynamic timeout: {ipc_timeout_seconds}s (2x baseline of {baseline_time_ms}ms)") + + logging.debug(f"pipe_name: {pipe_name}") + logging.debug(f"get_kern_arg_data args: {args}") + logging.debug(f"ipc_file_name: {ipc_file_name}") + + if not os.path.exists(pipe_name): + os.mkfifo(pipe_name) + os.chmod(pipe_name, stat.S_IRWXU | stat.S_IRWXG | stat.S_IRWXO) + + start_time = time.time() + with open(pipe_name, "rb") as fifo: # noqa: F841 + while True: + if time.time() - start_time > ipc_timeout_seconds: + raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data") + + try: + ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) + break + except Exception as e: + if time.time() - start_time > ipc_timeout_seconds: + raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data: {str(e)}") + time.sleep(0.1) + + type_map = { + "double*": ctypes.c_double, + "float*": ctypes.c_float, + "int*": ctypes.c_int, + "std::size_t*": ctypes.c_size_t, + "__half*": np.float16, + "__hip_bfloat16*": ml_dtypes.bfloat16, + } + + results = [] + pointer_args = list(filter(lambda arg: "*" in arg and "const" not in arg, args)) + logging.debug(f"pointer_args: {pointer_args}") + + for handle, arg, array_size in zip(ipc_handles, pointer_args, ptr_sizes): + ptr = open_ipc_handle(handle) + logging.debug(f"Opened IPC Ptr: {ptr} (0x{ptr:x})") + arg_type = arg.split()[0] + logging.debug(f"arg_type: {arg_type}") + + if arg_type in type_map: + dtype = type_map[arg_type] + logging.debug(f"dtype: {dtype}") + + # Special handling for FP16 and bfloat16 + if arg_type == "__half*": + temp_array = memcpy_d2h(ptr, array_size // 2, ctypes.c_uint16) + host_array = np.frombuffer(temp_array, dtype=np.float16) + elif arg_type == "__hip_bfloat16*": + temp_array = memcpy_d2h(ptr, array_size // 2, ctypes.c_uint16) + host_array = np.frombuffer(temp_array, dtype=ml_dtypes.bfloat16) + else: + num_elements = array_size // ctypes.sizeof(dtype) + host_array = memcpy_d2h(ptr, num_elements, dtype) + else: + raise TypeError(f"Unsupported pointer type: {arg_type}") + + logging.debug(f"Received data from IPC ({arg_type}/{len(host_array)}): {host_array}") + results.append(host_array) + + return results + diff --git a/src/accordo/config.py b/src/accordo/config.py new file mode 100644 index 00000000..b3c271d3 --- /dev/null +++ b/src/accordo/config.py @@ -0,0 +1,120 @@ +"""Configuration classes for Accordo validation.""" + +from dataclasses import dataclass, field +from typing import Union + + +@dataclass +class KernelArg: + """Represents a kernel argument with semantic information. + + Args: + name: Argument name (e.g., "result", "input", "count") + type: C/C++ type string (e.g., "double*", "const float*", "int") + direction: Data flow direction - "in", "out", "inout", or "auto" (infers from const) + + Examples: + >>> KernelArg(name="result", type="double*", direction="out") + >>> KernelArg(name="input", type="const double*", direction="in") + >>> KernelArg(name="count", type="unsigned long") # direction="auto" + """ + + name: str + type: str + direction: str = "auto" + + def __post_init__(self): + """Validate direction and infer from type if auto.""" + valid_directions = {"in", "out", "inout", "auto"} + if self.direction not in valid_directions: + raise ValueError(f"direction must be one of {valid_directions}, got '{self.direction}'") + + # Auto-infer direction from type if not specified + if self.direction == "auto": + if "const" in self.type and "*" in self.type: + self.direction = "in" + elif "*" in self.type: + self.direction = "out" # Default assumption for pointers + else: + self.direction = "in" # Scalars are typically inputs + + @classmethod + def from_string(cls, type_str: str, name: str = None) -> "KernelArg": + """Create KernelArg from a plain type string (backward compatibility). + + Args: + type_str: C/C++ type string + name: Optional argument name (auto-generated if not provided) + + Returns: + KernelArg instance + """ + if name is None: + name = f"arg_{id(type_str)}" # Generate unique name + return cls(name=name, type=type_str) + + @classmethod + def from_dict(cls, d: dict) -> "KernelArg": + """Create KernelArg from a dictionary.""" + return cls(**d) + + +@dataclass +class ValidationConfig: + """Configuration for Accordo kernel validation. + + Args: + kernel_name: Name of the kernel to validate + kernel_args: List of kernel arguments (KernelArg, str, or dict) + additional_includes: C++ include directives for custom types + tolerance: Absolute tolerance for array comparison + timeout_multiplier: Timeout = baseline_time_ms * timeout_multiplier + log_level: Logging level ("DEBUG", "INFO", "WARNING", "ERROR") + + Examples: + >>> config = ValidationConfig( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... KernelArg(name="result", type="double*", direction="out"), + ... KernelArg(name="input", type="const double*", direction="in"), + ... "int" # Plain string (backward compat) + ... ], + ... additional_includes=['"my_types.h"', ''], + ... tolerance=1e-6 + ... ) + """ + + kernel_name: str + kernel_args: list[Union[KernelArg, str, dict]] + additional_includes: list[str] = field(default_factory=list) + tolerance: float = 1e-6 + timeout_multiplier: float = 2.0 + log_level: str = "WARNING" + + def __post_init__(self): + """Normalize kernel_args to KernelArg instances.""" + normalized_args = [] + for i, arg in enumerate(self.kernel_args): + if isinstance(arg, KernelArg): + normalized_args.append(arg) + elif isinstance(arg, str): + # Convert plain string to KernelArg + normalized_args.append(KernelArg.from_string(arg, name=f"arg{i}")) + elif isinstance(arg, dict): + # Convert dict to KernelArg + if "name" not in arg: + arg["name"] = f"arg{i}" + normalized_args.append(KernelArg.from_dict(arg)) + else: + raise TypeError(f"kernel_args must be KernelArg, str, or dict, got {type(arg)}") + + self.kernel_args = normalized_args + + def get_arg_types(self) -> list[str]: + """Get list of argument type strings (for backward compatibility).""" + return [arg.type for arg in self.kernel_args] + + def get_arg_names(self) -> list[str]: + """Get list of argument names.""" + return [arg.name for arg in self.kernel_args] + diff --git a/src/accordo/exceptions.py b/src/accordo/exceptions.py new file mode 100644 index 00000000..f0f739e6 --- /dev/null +++ b/src/accordo/exceptions.py @@ -0,0 +1,36 @@ +"""Custom exceptions for Accordo validation.""" + + +class AccordoError(Exception): + """Base exception for all Accordo errors.""" + + pass + + +class AccordoBuildError(AccordoError): + """Raised when Accordo C++ library fails to build.""" + + pass + + +class AccordoTimeoutError(AccordoError): + """Raised when a kernel execution exceeds the timeout.""" + + def __init__(self, message: str, timeout_seconds: float): + super().__init__(message) + self.timeout_seconds = timeout_seconds + + +class AccordoProcessError(AccordoError): + """Raised when the instrumented process crashes or fails.""" + + def __init__(self, message: str, exit_code: int = None): + super().__init__(message) + self.exit_code = exit_code + + +class AccordoValidationError(AccordoError): + """Raised when array validation fails.""" + + pass + diff --git a/src/accordo/result.py b/src/accordo/result.py new file mode 100644 index 00000000..5664acda --- /dev/null +++ b/src/accordo/result.py @@ -0,0 +1,101 @@ +"""Result classes for Accordo validation.""" + +from dataclasses import dataclass +from typing import Optional + +import numpy as np + + +@dataclass +class ArrayMismatch: + """Represents a mismatch between reference and optimized arrays. + + Args: + arg_index: Index of the argument that failed validation + arg_name: Name of the argument + arg_type: Type string of the argument + max_difference: Maximum absolute difference between arrays + mean_difference: Mean absolute difference between arrays + reference_sample: Sample values from reference array + optimized_sample: Sample values from optimized array + """ + + arg_index: int + arg_name: str + arg_type: str + max_difference: float + mean_difference: float + reference_sample: np.ndarray + optimized_sample: np.ndarray + + def __str__(self) -> str: + """Human-readable string representation.""" + return ( + f"Mismatch in arg '{self.arg_name}' ({self.arg_type}): " + f"max_diff={self.max_difference:.2e}, mean_diff={self.mean_difference:.2e}" + ) + + +@dataclass +class ValidationResult: + """Result of Accordo validation. + + Args: + is_valid: True if all arrays matched within tolerance + error_message: Error message if validation failed + mismatches: List of array mismatches + matched_arrays: Dictionary of successfully matched arrays + execution_time_ms: Execution times for reference and optimized kernels + timeout_used: Timeout value used (if applicable) + """ + + is_valid: bool + error_message: Optional[str] = None + mismatches: list[ArrayMismatch] = None + matched_arrays: dict[str, dict] = None + execution_time_ms: dict[str, float] = None + timeout_used: Optional[float] = None + + def __post_init__(self): + """Initialize default values.""" + if self.mismatches is None: + self.mismatches = [] + if self.matched_arrays is None: + self.matched_arrays = {} + if self.execution_time_ms is None: + self.execution_time_ms = {} + + @property + def num_arrays_validated(self) -> int: + """Total number of arrays validated (matched + mismatched).""" + return len(self.matched_arrays) + len(self.mismatches) + + @property + def num_mismatches(self) -> int: + """Number of array mismatches.""" + return len(self.mismatches) + + @property + def success_rate(self) -> float: + """Percentage of arrays that matched.""" + total = self.num_arrays_validated + if total == 0: + return 0.0 + return (len(self.matched_arrays) / total) * 100.0 + + def summary(self) -> str: + """Get a human-readable summary of validation results.""" + if self.is_valid: + return f"✓ Validation passed! {self.num_arrays_validated} arrays matched within tolerance." + else: + lines = [f"✗ Validation failed: {self.error_message}"] + if self.mismatches: + lines.append(f"\nMismatched arrays ({len(self.mismatches)}):") + for mismatch in self.mismatches: + lines.append(f" - {mismatch}") + return "\n".join(lines) + + def __str__(self) -> str: + """String representation.""" + return self.summary() + From 4b4f3af58c3540eccdc1d8cff97c466e6a75e540 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 01:46:22 -0600 Subject: [PATCH 03/15] Intorduce take snapshot --- src/accordo/__init__.py | 46 ++- src/accordo/config.py | 27 +- src/accordo/snapshot.py | 58 +++ src/accordo/validator.py | 454 +++++++++++++++++++++++ src/intelliperf/formulas/formula_base.py | 160 ++++---- 5 files changed, 629 insertions(+), 116 deletions(-) create mode 100644 src/accordo/snapshot.py create mode 100644 src/accordo/validator.py diff --git a/src/accordo/__init__.py b/src/accordo/__init__.py index bc0df9d1..a1613ccd 100644 --- a/src/accordo/__init__.py +++ b/src/accordo/__init__.py @@ -3,21 +3,47 @@ Public API: - ValidationConfig: Configuration for kernel validation - KernelArg: Structured kernel argument representation + - Snapshot: Captured kernel argument data from binary execution - ValidationResult: Result of validation with detailed metrics - ArrayMismatch: Information about array validation failures - - AccordoValidator: Main validator class (TODO: implement) + - AccordoValidator: Main validator class for kernel validation - Exceptions: AccordoError, AccordoBuildError, AccordoTimeoutError, etc. -Example: - >>> from accordo import ValidationConfig, KernelArg +Quick Example (one-off validation): + >>> from accordo import ValidationConfig, KernelArg, AccordoValidator >>> config = ValidationConfig( ... kernel_name="my_kernel", ... kernel_args=[ - ... KernelArg(name="result", type="double*", direction="out"), - ... KernelArg(name="input", type="const double*", direction="in"), + ... KernelArg(name="result", type="double*"), + ... KernelArg(name="input", type="const double*"), ... ], ... tolerance=1e-6 ... ) + >>> validator = AccordoValidator(config) + >>> result = validator.validate( + ... reference_binary=["./app_ref"], + ... optimized_binary=["./app_opt"], + ... working_directory=".", + ... baseline_time_ms=10.0 + ... ) + +Efficient Example (multiple optimizations vs same reference): + >>> # Capture reference once (returns Snapshot object) + >>> ref_snapshot = validator.capture_snapshot( + ... binary=["./app_ref"], + ... working_directory=".", + ... timeout_seconds=30 + ... ) + >>> print(ref_snapshot) # Snapshot(binary='./app_ref', arrays=3, execution_time_ms=12.50) + >>> + >>> # Compare multiple optimizations + >>> for opt_binary in optimized_binaries: + ... opt_snapshot = validator.capture_snapshot( + ... binary=opt_binary, + ... working_directory=".", + ... timeout_seconds=60 + ... ) + ... result = validator.compare_snapshots(ref_snapshot, opt_snapshot) """ # Public API exports @@ -30,15 +56,21 @@ AccordoValidationError, ) from .result import ArrayMismatch, ValidationResult +from .snapshot import Snapshot +from .validator import AccordoValidator # Version -__version__ = "0.2.0-dev" +__version__ = "0.2.0" # Public API __all__ = [ # Config "ValidationConfig", "KernelArg", + # Validator + "AccordoValidator", + # Snapshot + "Snapshot", # Results "ValidationResult", "ArrayMismatch", @@ -48,6 +80,4 @@ "AccordoTimeoutError", "AccordoProcessError", "AccordoValidationError", - # TODO: Add AccordoValidator when implemented - # "AccordoValidator", ] diff --git a/src/accordo/config.py b/src/accordo/config.py index b3c271d3..6bcb5278 100644 --- a/src/accordo/config.py +++ b/src/accordo/config.py @@ -11,32 +11,19 @@ class KernelArg: Args: name: Argument name (e.g., "result", "input", "count") type: C/C++ type string (e.g., "double*", "const float*", "int") - direction: Data flow direction - "in", "out", "inout", or "auto" (infers from const) Examples: - >>> KernelArg(name="result", type="double*", direction="out") - >>> KernelArg(name="input", type="const double*", direction="in") - >>> KernelArg(name="count", type="unsigned long") # direction="auto" + >>> KernelArg(name="result", type="double*") + >>> KernelArg(name="input", type="const double*") + >>> KernelArg(name="count", type="unsigned long") + + Note: + Output arguments are identified by checking for "*" without "const" in the type. + This matches the existing IPC logic. """ name: str type: str - direction: str = "auto" - - def __post_init__(self): - """Validate direction and infer from type if auto.""" - valid_directions = {"in", "out", "inout", "auto"} - if self.direction not in valid_directions: - raise ValueError(f"direction must be one of {valid_directions}, got '{self.direction}'") - - # Auto-infer direction from type if not specified - if self.direction == "auto": - if "const" in self.type and "*" in self.type: - self.direction = "in" - elif "*" in self.type: - self.direction = "out" # Default assumption for pointers - else: - self.direction = "in" # Scalars are typically inputs @classmethod def from_string(cls, type_str: str, name: str = None) -> "KernelArg": diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py new file mode 100644 index 00000000..e496dba0 --- /dev/null +++ b/src/accordo/snapshot.py @@ -0,0 +1,58 @@ +"""Snapshot: Represents captured kernel argument data from a binary execution.""" + +from dataclasses import dataclass +from typing import List + +import numpy as np + + +@dataclass +class Snapshot: + """Represents a captured snapshot of kernel argument data. + + Attributes: + arrays: List of numpy arrays containing kernel argument data + 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 + + 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" + ... ) + >>> print(f"Captured {len(snapshot.arrays)} arrays in {snapshot.execution_time_ms}ms") + """ + + arrays: List[np.ndarray] + execution_time_ms: float + binary: List[str] + working_directory: str + + def __repr__(self) -> str: + """Pretty representation of snapshot.""" + binary_str = " ".join(self.binary) + return ( + f"Snapshot(binary='{binary_str}', " + f"arrays={len(self.arrays)}, " + f"execution_time_ms={self.execution_time_ms:.2f})" + ) + + def summary(self) -> str: + """Get a detailed summary of the snapshot.""" + binary_str = " ".join(self.binary) + lines = [ + f"Snapshot Summary:", + f" Binary: {binary_str}", + f" Working Directory: {self.working_directory}", + f" Execution Time: {self.execution_time_ms:.2f}ms", + f" Number of Arrays: {len(self.arrays)}", + ] + + for i, arr in enumerate(self.arrays): + lines.append(f" Array {i}: shape={arr.shape}, dtype={arr.dtype}") + + return "\n".join(lines) + diff --git a/src/accordo/validator.py b/src/accordo/validator.py new file mode 100644 index 00000000..94827948 --- /dev/null +++ b/src/accordo/validator.py @@ -0,0 +1,454 @@ +"""AccordoValidator: Main validation class for Accordo.""" + +import logging +import os +import signal +import subprocess +import time +from pathlib import Path +from typing import Optional + +import numpy as np + +from .config import ValidationConfig +from .exceptions import AccordoBuildError, AccordoProcessError, AccordoTimeoutError, AccordoValidationError +from .result import ArrayMismatch, ValidationResult +from .snapshot import Snapshot +from ._internal.codegen import generate_kernel_header +from ._internal.ipc.communication import get_kern_arg_data, send_response + + +class _TimeoutException(Exception): + """Internal exception for timeout handling.""" + pass + + +def _timeout_handler(signum, frame): + """Signal handler for timeout.""" + raise _TimeoutException("Operation timed out") + + +def _build_accordo(accordo_path: Path, parallel_jobs: int = 16) -> Path: + """Build Accordo C++ library. + + Args: + accordo_path: Path to Accordo directory + parallel_jobs: Number of parallel build jobs + + Returns: + Path to built library + + Raises: + AccordoBuildError: If build fails + """ + try: + # Configure with CMake + result = subprocess.run( + ["cmake", "-B", "build"], + cwd=accordo_path, + capture_output=True, + text=True, + check=True, + ) + logging.debug(f"CMake configure output: {result.stdout}") + + # Build + result = subprocess.run( + ["cmake", "--build", "build", "--parallel", str(parallel_jobs)], + cwd=accordo_path, + capture_output=True, + text=True, + check=True, + ) + logging.debug(f"CMake build output: {result.stdout}") + + lib_path = accordo_path / "build" / "lib" / "libaccordo.so" + if not lib_path.exists(): + raise AccordoBuildError(f"Library not found at {lib_path}") + + return lib_path + + except subprocess.CalledProcessError as e: + raise AccordoBuildError(f"Accordo build failed: {e.stderr}") + except Exception as e: + raise AccordoBuildError(f"Accordo build failed: {str(e)}") + + +def _validate_arrays(arr1: np.ndarray, arr2: np.ndarray, tolerance: float) -> bool: + """Validate two arrays are close within tolerance. + + Args: + arr1: First array + arr2: Second array + tolerance: Absolute tolerance + + Returns: + True if arrays match within tolerance + """ + return np.allclose(arr1, arr2, atol=tolerance, rtol=0) + + +class AccordoValidator: + """Validator for GPU kernel correctness using Accordo. + + This class manages the entire validation pipeline: + - Building the Accordo C++ library + - Running instrumented processes + - Collecting kernel argument data via IPC + - Validating arrays match within tolerance + + Example: + >>> from accordo import AccordoValidator, ValidationConfig, KernelArg + >>> config = ValidationConfig( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... KernelArg(name="result", type="double*", direction="out"), + ... KernelArg(name="input", type="const double*", direction="in"), + ... ], + ... tolerance=1e-6 + ... ) + >>> validator = AccordoValidator(config) + >>> result = validator.validate(reference_app, optimized_app) + >>> if result.is_valid: + ... print("Validation passed!") + """ + + def __init__( + self, + config: ValidationConfig, + accordo_path: Optional[Path] = None, + force_rebuild: bool = False, + parallel_jobs: int = 16, + ): + """Initialize AccordoValidator. + + Args: + config: Validation configuration + accordo_path: Path to Accordo directory (auto-detected if None) + force_rebuild: Force rebuild even if library exists + parallel_jobs: Number of parallel build jobs + """ + self.config = config + self.parallel_jobs = parallel_jobs + self._built = False + self._lib_path = None + + # Auto-detect accordo_path if not provided + if accordo_path is None: + # Try to find it relative to this file + accordo_dir = Path(__file__).parent + if (accordo_dir / "build").exists() or (accordo_dir / "CMakeLists.txt").exists(): + accordo_path = accordo_dir + else: + # Try environment variable + from intelliperf.utils.env import get_accordo_path + + accordo_path = Path(get_accordo_path()) + + self.accordo_path = Path(accordo_path) + logging.debug(f"Accordo path: {self.accordo_path}") + + # Build if forced or library doesn't exist + lib_path = self.accordo_path / "build" / "lib" / "libaccordo.so" + if force_rebuild or not lib_path.exists(): + logging.info("Building Accordo C++ library...") + self._lib_path = _build_accordo(self.accordo_path, parallel_jobs) + self._built = True + else: + self._lib_path = lib_path + self._built = True + + def capture_snapshot( + self, + binary: list[str], + working_directory: str = ".", + timeout_seconds: int = 30, + ) -> Snapshot: + """Capture a snapshot of kernel argument data from a binary execution. + + Args: + binary: Command to run binary (e.g., ["./app", "arg1"]) + working_directory: Directory to run binary from + timeout_seconds: Timeout for this capture + + Returns: + Snapshot object containing captured arrays and execution metadata + + Raises: + AccordoBuildError: If Accordo library not built + AccordoProcessError: If instrumented process crashes + AccordoTimeoutError: If execution exceeds timeout + + Note: + Binary must be pre-compiled. Accordo does not build applications. + """ + if not self._built: + raise AccordoBuildError("Accordo library not built") + + # Generate kernel header with additional includes + arg_types = self.config.get_arg_types() + generate_kernel_header(arg_types, self.config.additional_includes) + + # Wrap app run with timeout + old_handler = signal.signal(signal.SIGALRM, _timeout_handler) + signal.alarm(timeout_seconds) + + try: + start_time = time.time() + result_arrays = self._run_instrumented_app( + binary, + working_directory, + label="snapshot", + baseline_time_ms=None + ) + signal.alarm(0) # Cancel alarm on success + execution_time_ms = (time.time() - start_time) * 1000 + + return Snapshot( + arrays=result_arrays, + execution_time_ms=execution_time_ms, + binary=binary, + working_directory=working_directory, + ) + except _TimeoutException: + signal.alarm(0) + logging.error(f"Snapshot capture timed out after {timeout_seconds}s") + raise AccordoTimeoutError( + f"Snapshot capture timed out after {timeout_seconds}s. " + "This may indicate a GPU crash or hung process.", + timeout_seconds=timeout_seconds + ) + except TimeoutError as e: + signal.alarm(0) + raise AccordoTimeoutError(f"Snapshot timeout: {str(e)}", timeout_seconds) + except RuntimeError as e: + signal.alarm(0) + raise AccordoProcessError(f"Process crashed during snapshot: {str(e)}") + finally: + signal.alarm(0) + signal.signal(signal.SIGALRM, old_handler) + + def compare_snapshots( + self, + reference_snapshot: Snapshot, + optimized_snapshot: Snapshot, + ) -> ValidationResult: + """Compare two snapshots and validate their arrays. + + Args: + reference_snapshot: Snapshot from reference binary (from capture_snapshot) + optimized_snapshot: Snapshot from optimized binary (from capture_snapshot) + + Returns: + ValidationResult with validation status and details + """ + results = { + "reference": reference_snapshot.arrays, + "optimized": optimized_snapshot.arrays, + } + execution_times = { + "reference": reference_snapshot.execution_time_ms, + "optimized": optimized_snapshot.execution_time_ms, + } + + return self._validate_results(results, execution_times) + + def validate( + self, + reference_binary: list[str], + optimized_binary: list[str], + working_directory: str = ".", + baseline_time_ms: Optional[float] = None, + ) -> ValidationResult: + """Validate optimized kernel against reference (convenience method). + + This is a convenience wrapper that captures both snapshots and compares them. + For better performance when validating multiple optimizations against the same + reference, use capture_snapshot() and compare_snapshots() directly. + + Args: + reference_binary: Command to run reference binary (e.g., ["./app", "arg1"]) + optimized_binary: Command to run optimized binary (e.g., ["./app_opt", "arg1"]) + working_directory: Directory to run binaries from + baseline_time_ms: Baseline execution time for dynamic timeout + + Returns: + ValidationResult with validation status and details + + Raises: + AccordoBuildError: If Accordo library not built + AccordoProcessError: If instrumented process crashes + AccordoTimeoutError: If execution exceeds timeout + + Note: + Both binaries must be pre-compiled. Accordo does not build applications. + """ + # Calculate timeouts + ref_timeout = 30 # Default for reference + if baseline_time_ms: + opt_timeout = int((baseline_time_ms * self.config.timeout_multiplier / 1000.0) + 30.0) + else: + opt_timeout = 30 + + # Capture snapshots + reference_snapshot = self.capture_snapshot(reference_binary, working_directory, ref_timeout) + optimized_snapshot = self.capture_snapshot(optimized_binary, working_directory, opt_timeout) + + # Compare + 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 + ) -> list[np.ndarray]: + """Run an instrumented application and collect kernel argument data. + + Args: + binary_cmd: Binary command with arguments (e.g., ["./app", "arg1"]) + working_directory: Directory to run the binary from + label: Label for this run ("reference" or "optimized") + baseline_time_ms: Baseline time for dynamic timeout + + Returns: + List of numpy arrays with kernel argument data + """ + timestamp = int(time.time() * 1000) # Use milliseconds for uniqueness + pipe_name = f"/tmp/kernel_pipe_{timestamp}_{label}" + ipc_file_name = f"/tmp/ipc_handle_{timestamp}_{label}.bin" + + # Clean up any existing files + for file_path in [pipe_name, ipc_file_name]: + if os.path.exists(file_path): + os.remove(file_path) + + # Set up environment + env = os.environ.copy() + env["HSA_TOOLS_LIB"] = str(self._lib_path) + env["KERNEL_TO_TRACE"] = self.config.kernel_name + + # Set log level + debug_level = logging.getLogger().getEffectiveLevel() + level_map = { + logging.WARNING: 0, + logging.INFO: 1, + logging.DEBUG: 2, + logging.NOTSET: 3, + } + env["ACCORDO_LOG_LEVEL"] = str(level_map.get(debug_level, 0)) + env["ACCORDO_PIPE_NAME"] = pipe_name + env["ACCORDO_IPC_OUTPUT_FILE"] = ipc_file_name + + # Launch process + logging.debug(f"Launching {label} process with PID for kernel {self.config.kernel_name}") + logging.debug(f"binary_cmd: {binary_cmd}") + logging.debug(f"working_directory: {working_directory}") + logging.debug(f"kernel_args: {self.config.get_arg_types()}") + logging.debug(f"ipc_file_name: {ipc_file_name}") + + original_dir = os.getcwd() + try: + os.chdir(working_directory) + process_pid = os.posix_spawn(binary_cmd[0], binary_cmd, env) + logging.debug(f"Launched {label} process with PID: {process_pid}") + finally: + os.chdir(original_dir) + + # Get kernel argument data via IPC + try: + result_arrays = get_kern_arg_data( + pipe_name, + self.config.get_arg_types(), + ipc_file_name, + process_pid=process_pid, + baseline_time_ms=baseline_time_ms, + ) + except TimeoutError as e: + # Kill the process if it timed out + try: + os.kill(process_pid, 9) + except: + pass + raise + + # Send completion response + send_response(pipe_name) + + return result_arrays + + def _validate_results( + self, results: dict[str, list[np.ndarray]], execution_times: dict[str, float] + ) -> ValidationResult: + """Validate results from reference and optimized runs. + + Args: + results: Dictionary with "reference" and "optimized" array lists + execution_times: Execution times for each run + + Returns: + ValidationResult with validation status + """ + reference_arrays = results["reference"] + optimized_arrays = results["optimized"] + + if len(reference_arrays) != len(optimized_arrays): + return ValidationResult( + is_valid=False, + error_message=f"Array count mismatch: {len(reference_arrays)} vs {len(optimized_arrays)}", + execution_time_ms=execution_times, + ) + + mismatches = [] + matched_arrays = {} + + for i, (ref_arr, opt_arr) in enumerate(zip(reference_arrays, optimized_arrays)): + arg = self.config.kernel_args[i] + + if not _validate_arrays(ref_arr, opt_arr, self.config.tolerance): + # Array mismatch + diff = np.abs(ref_arr - opt_arr) + mismatch = ArrayMismatch( + arg_index=i, + arg_name=arg.name, + arg_type=arg.type, + max_difference=float(np.max(diff)), + mean_difference=float(np.mean(diff)), + reference_sample=ref_arr[:10] if len(ref_arr) > 10 else ref_arr, + optimized_sample=opt_arr[:10] if len(opt_arr) > 10 else opt_arr, + ) + mismatches.append(mismatch) + + logging.debug(f"Arrays at index {i} for arg '{arg.name}' ({arg.type}) are NOT close.") + logging.debug(f" Max difference: {mismatch.max_difference}") + logging.debug(f" Mean difference: {mismatch.mean_difference}") + else: + # Array matched + matched_arrays[arg.name] = { + "index": i, + "type": arg.type, + "size": len(ref_arr), + } + logging.debug(f"Arrays at index {i} for arg '{arg.name}' ({arg.type}) are close.") + + # Determine overall success + is_valid = len(mismatches) == 0 + + if is_valid: + return ValidationResult( + is_valid=True, + matched_arrays=matched_arrays, + execution_time_ms=execution_times, + ) + else: + # Build error message + error_lines = [f"Validation failed: {len(mismatches)} array(s) mismatched"] + for m in mismatches: + error_lines.append(f" - {m}") + error_message = "\n".join(error_lines) + + return ValidationResult( + is_valid=False, + error_message=error_message, + mismatches=mismatches, + matched_arrays=matched_arrays, + execution_time_ms=execution_times, + ) + diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 35001319..88d1b190 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -37,9 +37,7 @@ import numpy as np import pandas as pd -from accordo.python.code_gen import generate_header -from accordo.python.communicate import get_kern_arg_data, send_response -from accordo.python.utils import run_subprocess +from accordo import AccordoValidator, KernelArg, ValidationConfig from intelliperf import __version__ from intelliperf.core.application import Application from intelliperf.core.logger import Logger @@ -237,6 +235,10 @@ def __init__( # Store num_attempts self.num_attempts = num_attempts + # Accordo caching (for efficient multi-iteration validation) + self._accordo_validator = None + self._reference_snapshot = None + # Initialize logger self._logger = Logger(run_name=name) self._logger.record( @@ -420,7 +422,10 @@ def optimize_pass(self, target_kernel: str = None): @abstractmethod def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tolerance: float = 1e-6): """ - Validates the the application. + Validates the application using Accordo. + + Uses snapshot caching: reference app is captured once on first call, + then each optimized version is captured and compared to the cached reference. """ if self.unittest_command: success, output = self._application.run_unit_test() @@ -431,96 +436,75 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole ) return Result(success=True, asset={"log": output}) + # Build the optimized application first (Accordo doesn't build) self._application.build() - unoptimized_binary = self._application.get_app_cmd()[0] - optimized_binary = self._reference_app.get_app_cmd()[0] - - logging.debug(f"unoptimized_binary: {unoptimized_binary}") - logging.debug(f"optimized_binary: {optimized_binary}") - - accordo_directory = get_accordo_path() - - results = {} - for app, label in zip([self._reference_app, self._application], ["unoptimized", "optimized"]): - logging.debug(f"Running accordo for {label}") - timestamp = int(time.time()) - pipe_name = f"/tmp/kernel_pipe_{timestamp}" - ipc_file_name = f"/tmp/ipc_handle_{timestamp}.bin" - - for file in [ipc_file_name, ipc_file_name]: - if os.path.exists(file): - os.remove(file) - generate_header(kernel_args) - - run_subprocess(["cmake", "-B", "build"], accordo_directory) - run_subprocess(["cmake", "--build", "build", "--parallel", "16"], accordo_directory) - lib = os.path.join(accordo_directory, "build", "lib", "libaccordo.so") - env = os.environ.copy() - env["HSA_TOOLS_LIB"] = lib - env["KERNEL_TO_TRACE"] = kernel - - # Get the debug level from logger and convert it - debug_level = logging.getLogger().getEffectiveLevel() - level_map = { - logging.WARNING: 0, # Warning - logging.INFO: 1, # Info - logging.DEBUG: 2, # Debug - logging.NOTSET: 3, # NOTEST - } - env["ACCORDO_LOG_LEVEL"] = str(level_map.get(debug_level, 0)) # Default to 0 (Warning) if level not found - env["ACCORDO_PIPE_NAME"] = pipe_name - env["ACCORDO_IPC_OUTPUT_FILE"] = ipc_file_name - - binary = app.get_app_cmd_without_args() - binary_with_args = app.get_app_cmd() - project_directory = app.get_project_directory() - logging.debug(f"binary: {binary}") - logging.debug(f"project_directory: {project_directory}") - logging.debug(f"kernel: {kernel}") - logging.debug(f"binary_with_args: {binary_with_args}") - logging.debug(f"kernel_args: {kernel_args}") - logging.debug(f"ipc_file_name: {ipc_file_name}") - - original_dir = os.getcwd() - os.chdir(project_directory) - os.posix_spawn(binary, binary_with_args, env) - os.chdir(original_dir) + # Create validator if not already created (and cache it) + if self._accordo_validator is None: + kernel_arg_objects = [KernelArg(name=f"arg{i}", type=arg_type) for i, arg_type in enumerate(kernel_args)] + + config = ValidationConfig( + kernel_name=kernel, + kernel_args=kernel_arg_objects, + tolerance=accordo_absolute_tolerance, + timeout_multiplier=2.0, + ) + + self._accordo_validator = AccordoValidator(config) + logging.debug("Created and cached Accordo validator") + + # Capture reference snapshot if not already captured (and cache it) + if self._reference_snapshot is None: try: - results[label] = get_kern_arg_data(pipe_name, kernel_args, ipc_file_name) - except TimeoutError as e: - logging.error(f"Timeout while getting kernel argument data for {label}: {str(e)}") - return Result( - success=False, - error_report=f"Timeout while getting kernel argument data for {label}: {str(e)}. The code may have crashed.", + reference_binary = self._reference_app.get_app_cmd() + working_dir = self._reference_app.get_project_directory() + + logging.debug("Capturing reference snapshot (will be cached)") + self._reference_snapshot = self._accordo_validator.capture_snapshot( + binary=reference_binary, + working_directory=working_dir, + timeout_seconds=30 ) - send_response(pipe_name) - logging.debug(f"results unoptimized: {results['unoptimized']}") - logging.debug(f"results optimized: {results['optimized']}") - key0, key1 = results.keys() - for i in range(len(results[key0])): - if not validate_arrays(results[key0][i], results[key1][i], accordo_absolute_tolerance): - diff = np.abs(results[key0][i] - results[key1][i]) - logging.debug(f"Arrays at index {i} for '{key0}' and '{key1}' are NOT close.") - logging.debug(f" {key0}[{i}]: {results[key0][i]}") - logging.debug(f" {key1}[{i}]: {results[key1][i]}") - logging.debug(f" Difference: {diff}") - logging.debug(f" Max difference: {np.max(diff)}") + logging.debug(f"Reference snapshot captured in {self._reference_snapshot.execution_time_ms:.2f}ms") + except Exception as e: + logging.error(f"Failed to capture reference snapshot: {str(e)}") + return Result(success=False, error_report=f"Failed to capture reference snapshot: {str(e)}") + + # Capture optimized snapshot and compare with cached reference + try: + baseline_time = getattr(self, "baseline_time_ms", None) + optimized_binary = self._application.get_app_cmd() + working_dir = self._application.get_project_directory() + + # Calculate timeout for optimized + if baseline_time: + opt_timeout = int((baseline_time * 2.0 / 1000.0) + 30.0) + else: + opt_timeout = 30 + logging.debug("Capturing optimized snapshot") + optimized_snapshot = self._accordo_validator.capture_snapshot( + binary=optimized_binary, + working_directory=working_dir, + timeout_seconds=opt_timeout + ) + logging.debug(f"Optimized snapshot captured in {optimized_snapshot.execution_time_ms:.2f}ms") + + # Compare snapshots + validation_result = self._accordo_validator.compare_snapshots( + self._reference_snapshot, + optimized_snapshot + ) + + if validation_result.is_valid: + logging.debug("Validation succeeded.") + return Result(success=True) else: - argument_name = kernel_args[i] - logging.debug( - f"Arrays at index {i} for '{key0}' and '{key1}' are close. The argument type is '{argument_name}'." - ) - for i in range(len(results[key0])): - if not validate_arrays(results[key0][i], results[key1][i], accordo_absolute_tolerance): - argument_name = kernel_args[i] - return Result( - success=False, - error_report=f"The optimized code output does not match the unoptimized code output. Values at index {i} for the '{argument_name}' pointer are NOT close.", - ) - logging.debug("Validation succeeded.") - return Result(success=True) + return Result(success=False, error_report=validation_result.error_message) + + except Exception as e: + logging.error(f"Accordo validation error: {str(e)}") + return Result(success=False, error_report=f"Accordo validation error: {str(e)}") @abstractmethod def performance_validation_pass(self): From 51b0935674326528bf2932c957deb782df2c7d23 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 01:47:59 -0600 Subject: [PATCH 04/15] Remove logging files --- accordo.todo | 334 ------------------------------ src/accordo/README_REFACTORING.md | 168 --------------- 2 files changed, 502 deletions(-) delete mode 100644 accordo.todo delete mode 100644 src/accordo/README_REFACTORING.md diff --git a/accordo.todo b/accordo.todo deleted file mode 100644 index edbca5e4..00000000 --- a/accordo.todo +++ /dev/null @@ -1,334 +0,0 @@ -# Accordo API Refactoring TODO - -## Current Problems -- [ ] No clear abstraction layers - CMake builds, pipe management, IPC handling all exposed -- [ ] Scattered concerns - code_gen, communicate, utils in separate modules -- [ ] Hard to use - Formula authors must manage pipes, env vars, processes manually -- [ ] Poor testability - Everything coupled to formula_base -- [ ] No clean imports - `from accordo.python.code_gen import generate_header` is ugly -- [ ] Strings for kernel args - no semantic info, bad for LLM generation - -## Proposed Clean API Design - -### High-Level User API (3 Main Classes) - -```python -from accordo import AccordoValidator, ValidationConfig, ValidationResult - -# Clean, simple usage -config = ValidationConfig( - kernel_name="my_kernel", - kernel_args=[ - KernelArg(name="result", type="double*", direction="out"), - KernelArg(name="input", type="const double*", direction="in"), - KernelArg(name="count", type="unsigned long", direction="in") - ], - additional_includes=['"my_types.h"'], # For custom types - tolerance=1e-6, - timeout_multiplier=2.0 -) - -validator = AccordoValidator(config) -result = validator.validate(reference_app, optimized_app, baseline_time_ms=10.5) - -if result.is_valid: - print(f"✓ Validation passed! {result.num_arrays_validated} arrays matched") -else: - print(f"✗ Failed: {result.error_message}") - for m in result.mismatches: - print(f" Arg '{m.arg_name}' ({m.arg_type}): max_diff={m.max_difference}") -``` - -### Package Structure - -``` -src/accordo/ -├── __init__.py # Public API: AccordoValidator, ValidationConfig, ValidationResult -├── validator.py # AccordoValidator class -├── config.py # ValidationConfig, KernelArg dataclasses -├── result.py # ValidationResult, ArrayMismatch dataclasses -├── exceptions.py # AccordoError, AccordoBuildError, AccordoTimeoutError, etc. -├── _internal/ # Private implementation (not imported by users) -│ ├── __init__.py -│ ├── builder.py # AccordoBuilder - manages CMake builds -│ ├── runtime.py # AccordoRuntime - manages process + IPC -│ ├── codegen.py # generate_kernel_header() - creates KernelArguments.hpp -│ ├── ipc/ -│ │ ├── __init__.py -│ │ ├── communication.py # IPC handle reading/writing -│ │ ├── memory.py # Device-to-host memory operations -│ │ └── types.py # Type mappings (float*, int*, etc) -│ └── hip.py # HIP interop functions -└── cpp/ # C++ library (unchanged) - ├── accordo.hpp - ├── accordo.hip - └── CMakeLists.txt -``` - -## Key Design Decisions - -### 1. Structured KernelArg (Not Plain Strings) -**Rationale:** LLMs generate JSON better than parsing strings. Named args give better error messages. - -```python -# ✅ Structured (LLM-friendly, semantic, extensible) -kernel_args=[ - KernelArg(name="result", type="double*", direction="out"), - KernelArg(name="input", type="const double*", direction="in") -] - -# ❌ Old way (hard for LLM, unclear semantics) -kernel_args=["double*", "const double*"] -``` - -**Benefits:** -- LLMs excel at generating structured JSON -- Better error messages: "Arg 'result' failed" vs "Arg 0 failed" -- Explicit input/output direction -- Easy to extend with metadata (size, validate flag, etc) -- Backward compat: still accept strings, auto-convert to KernelArg - -### 2. No BuildConfig Class -**Rationale:** Build is deterministic, fast, and users don't care about build details. - -- Accordo auto-builds on first use -- CMake caching makes rebuilds instant -- Optional kwargs for power users: `force_rebuild=True`, `parallel_jobs=16` - -### 3. Flat Arguments Only (For Now) -**Rationale:** 95% of GPU kernels use flat signatures. Nested struct support can be added later. - -```python -# ✅ Supported (flat args) -kernel_args=[ - KernelArg("output", "float*"), - KernelArg("input", "const float*"), - KernelArg("size", "int") -] - -# ❌ Not supported yet (nested) -kernel_args=[ - KernelArg("matrix", "Matrix*") # Matrix has nested pointers -] - -# Future: Add visitor/accessor pattern for nested types -``` - -### 4. additional_includes for Custom Types -**Rationale:** Accordo generates C++ header, needs includes to define custom types. - -```python -config = ValidationConfig( - kernel_name="my_kernel", - kernel_args=[ - KernelArg("matrix", "MyMatrix*"), - KernelArg("data", "__hip_bfloat16*") - ], - additional_includes=[ - '"my_app/matrix.h"', # Custom types - '' # HIP types - ] -) -``` - -Generated header: -```cpp -#pragma once -#include -#include "my_app/matrix.h" // User's custom types -#include // HIP types - -struct KernelArguments { - MyMatrix* arg0; - __hip_bfloat16* arg1; - - auto as_tuple() const { - return std::tie(arg0, arg1); - } -}; -``` - -## Core Classes Detail - -### ValidationConfig -```python -@dataclass -class KernelArg: - name: str # "result", "input", "count" - type: str # "double*", "const float*", "int" - direction: str = "auto" # "in", "out", "inout", "auto" (infer from const) - -@dataclass -class ValidationConfig: - kernel_name: str - kernel_args: list[Union[KernelArg, str, dict]] # Accepts multiple formats - additional_includes: list[str] = field(default_factory=list) - tolerance: float = 1e-6 - timeout_multiplier: float = 2.0 - log_level: str = "WARNING" -``` - -### ValidationResult -```python -@dataclass -class ArrayMismatch: - arg_index: int - arg_name: str - arg_type: str - max_difference: float - mean_difference: float - reference_sample: np.ndarray - optimized_sample: np.ndarray - -@dataclass -class ValidationResult: - is_valid: bool - error_message: Optional[str] - mismatches: list[ArrayMismatch] - matched_arrays: dict[str, dict] - execution_time_ms: dict[str, float] # {"reference": 10.5, "optimized": 8.3} - timeout_used: Optional[float] - - @property - def num_arrays_validated(self) -> int: - return len(self.matched_arrays) + len(self.mismatches) -``` - -### AccordoValidator -```python -class AccordoValidator: - def __init__( - self, - config: ValidationConfig, - accordo_path: Optional[Path] = None, # Auto-detected - force_rebuild: bool = False, - parallel_jobs: int = 16 - ): - """Initialize validator. Lazy builds on first validate() call.""" - - def validate( - self, - reference_app: Application, - optimized_app: Application, - baseline_time_ms: Optional[float] = None - ) -> ValidationResult: - """ - Validate optimized vs reference. - - Args: - baseline_time_ms: For dynamic timeout (timeout = baseline * multiplier) - - Returns: - ValidationResult with validation status and details - - Raises: - AccordoTimeoutError, AccordoProcessError, AccordoBuildError - """ -``` - -### Exceptions -```python -class AccordoError(Exception): pass -class AccordoBuildError(AccordoError): pass -class AccordoTimeoutError(AccordoError): - def __init__(self, message, timeout_seconds): - self.timeout_seconds = timeout_seconds -class AccordoProcessError(AccordoError): - def __init__(self, message, exit_code=None): - self.exit_code = exit_code -``` - -## Integration with IntelliPerf - -### Simplified formula_base.py -```python -from accordo import AccordoValidator, ValidationConfig, KernelArg - -def correctness_validation_pass(self, kernel, kernel_args, tolerance=1e-6): - # Create config - config = ValidationConfig( - kernel_name=kernel, - kernel_args=kernel_args, # List of KernelArg or strings - tolerance=tolerance - ) - - # Validate - validator = AccordoValidator(config) - result = validator.validate( - reference_app=self._reference_app, - optimized_app=self._application, - baseline_time_ms=getattr(self, "baseline_time_ms", None) - ) - - return Result( - success=result.is_valid, - error_report=result.error_message if not result.is_valid else "" - ) -``` - -## Implementation Tasks - -- [ ] Create new package structure -- [ ] Implement config.py (ValidationConfig, KernelArg) -- [ ] Implement result.py (ValidationResult, ArrayMismatch) -- [ ] Implement exceptions.py (all custom exceptions) -- [ ] Implement validator.py (AccordoValidator main class) -- [ ] Implement _internal/builder.py (CMake build management) -- [ ] Implement _internal/runtime.py (process + IPC management) -- [ ] Implement _internal/codegen.py (header generation with additional_includes) -- [ ] Implement _internal/ipc/ (communication, memory, types) -- [ ] Update __init__.py (clean public API exports) -- [ ] Update formula_base.py to use new API -- [ ] Write tests for each component -- [ ] Update documentation -- [ ] Migration guide for existing code - -## Future Enhancements (Post-Flat-Args) - -### Nested Struct Support -```python -# Option 1: User provides accessor -kernel_args=[ - KernelArg( - name="matrix", - type="Matrix*", - nested_pointers=["Matrix.data"], # Specify nested pointers to validate - element_type="float*" - ) -] - -# Option 2: Type definitions for reflection -type_definitions={ - "Matrix": { - "fields": [ - {"name": "data", "type": "float*", "validate": True}, - {"name": "rows", "type": "int", "validate": False} - ] - } -} -``` - -### Size Expressions -```python -kernel_args=[ - KernelArg(name="output", type="float*", size_expr="N"), - KernelArg(name="input", type="float*", size_expr="N*M") -] -``` - -### Selective Validation -```python -kernel_args=[ - KernelArg(name="output", type="float*", validate=True), - KernelArg(name="config", type="Config*", validate=False) # Skip complex structs -] -``` - -## Notes -- Keep API minimal and clean - no BuildConfig -- Optimize for LLM generation - structured KernelArg over strings -- Start with flat args (95% use case), add nested later -- Auto-build, auto-detect, minimize user configuration -- Clear error messages with named arguments -- Type hints everywhere for IDE support - diff --git a/src/accordo/README_REFACTORING.md b/src/accordo/README_REFACTORING.md deleted file mode 100644 index a788e38e..00000000 --- a/src/accordo/README_REFACTORING.md +++ /dev/null @@ -1,168 +0,0 @@ -# Accordo API Refactoring Progress - -## ✅ Completed - -### 1. Package Structure -- Created `_internal/` directory for private implementation -- Created `_internal/ipc/` subdirectory for IPC modules -- Added proper `__init__.py` files - -### 2. Public API Classes - -#### `config.py` - Configuration -- ✅ `KernelArg`: Structured kernel argument with name, type, and direction - - Supports auto-inference of direction from type - - Backward compatible with plain strings and dicts -- ✅ `ValidationConfig`: Configuration for validation - - Supports `kernel_args` as KernelArg, str, or dict - - Includes `additional_includes` for custom types - - Configurable tolerance and timeout - -#### `result.py` - Results -- ✅ `ArrayMismatch`: Detailed mismatch information -- ✅ `ValidationResult`: Validation result with metrics - - Properties: `num_arrays_validated`, `success_rate` - - Method: `summary()` for human-readable output - -#### `exceptions.py` - Exceptions -- ✅ `AccordoError`: Base exception -- ✅ `AccordoBuildError`: Build failures -- ✅ `AccordoTimeoutError`: Timeout errors (with timeout_seconds attribute) -- ✅ `AccordoProcessError`: Process crashes (with exit_code attribute) -- ✅ `AccordoValidationError`: Validation failures - -### 3. Internal Modules - -#### `_internal/codegen.py` -- ✅ `generate_kernel_header()`: Creates KernelArguments.hpp - - Supports `additional_includes` parameter - - Generates flat argument structures - -#### `_internal/hip.py` -- ✅ `hip_try()`: HIP error checking -- ✅ `open_ipc_handle()`: Open IPC memory handles -- ✅ `memcpy_d2h()`: Device-to-host memory copy - -#### `_internal/ipc/communication.py` -- ✅ `read_ipc_handles()`: Read IPC handles from file -- ✅ `send_response()`: Send completion signal via pipe -- ✅ `get_kern_arg_data()`: Get kernel argument data via IPC - - Supports dynamic timeout based on baseline - - Type mapping for common GPU types (float*, double*, __half*, etc.) - -### 4. Public API Export -- ✅ Updated `__init__.py` to export public API classes -- ✅ Clean imports: `from accordo import ValidationConfig, KernelArg, ValidationResult` - -## 🚧 TODO - -### 1. AccordoValidator Implementation -The main validator class needs to be implemented. This will: -- Wrap the existing functionality from `formula_base.py` -- Manage Accordo C++ library building (via CMake) -- Handle process management and IPC communication -- Provide the `validate(reference_app, optimized_app)` method - -**Key Design Decisions:** -- Lazy build on first use -- Auto-detect Accordo path -- Dynamic timeout calculation -- Clean error handling with custom exceptions - -### 2. Builder Module (`_internal/builder.py`) -Manage CMake builds of Accordo C++ library: -- Check if build needed -- Run CMake with proper flags -- Handle build errors -- Cache build results - -### 3. Runtime Module (`_internal/runtime.py`) -Manage process execution and IPC: -- Launch instrumented processes -- Monitor process health -- Handle timeouts -- Clean up resources - -### 4. Integration with IntelliPerf -Update `formula_base.py` to use the new API: -```python -from accordo import AccordoValidator, ValidationConfig, KernelArg - -config = ValidationConfig( - kernel_name=kernel, - kernel_args=[KernelArg(name=f"arg{i}", type=t) for i, t in enumerate(kernel_args)], - tolerance=tolerance -) - -validator = AccordoValidator(config) -result = validator.validate( - reference_app=self._reference_app, - optimized_app=self._application, - baseline_time_ms=getattr(self, "baseline_time_ms", None) -) - -return Result( - success=result.is_valid, - error_report=result.error_message if not result.is_valid else "" -) -``` - -### 5. Tests -- Unit tests for each module -- Integration tests with real kernels -- Test backward compatibility - -### 6. Documentation -- API reference docs -- Usage examples -- Migration guide - -## 📝 Design Notes - -### Flat Arguments Only (For Now) -Current implementation supports only flat argument lists: -```python -# ✅ Supported -kernel_args=[ - KernelArg("output", "float*"), - KernelArg("input", "const float*"), - KernelArg("size", "int") -] - -# ❌ Not yet supported (nested structs) -kernel_args=[ - KernelArg("matrix", "Matrix*") # Matrix has nested pointers -] -``` - -Future enhancement: Add visitor/accessor pattern for nested types. - -### Backward Compatibility -The API accepts multiple input formats: -- `KernelArg` objects (new, preferred) -- Plain strings (backward compatible) -- Dicts (for JSON/LLM generation) - -All formats are automatically normalized to `KernelArg` instances. - -### Type Mapping -Supported GPU types in `_internal/ipc/communication.py`: -- `double*`, `float*` -- `int*`, `std::size_t*` -- `__half*` (FP16) -- `__hip_bfloat16*` (BFloat16) - -Additional types can be added to the `type_map` dictionary. - -## 🎯 Next Steps - -1. **Implement AccordoValidator** - This is the main missing piece -2. **Test with existing formulas** - Ensure backward compatibility -3. **Update formula_base.py** - Integrate new API -4. **Add comprehensive tests** - Cover edge cases -5. **Document migration path** - Help users adopt new API - -## 📚 Reference - -See `accordo.todo` for the complete design document. - From 96b624ca66f61ecf024e0666688175ec7f675058 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 02:05:25 -0600 Subject: [PATCH 05/15] Update examples --- .../uncoalesced/uncoalesced.hip | 6 +-- .../b2b_matrix_transpose.hip | 22 +++++++--- .../matrix_transpose/matrix_transpose.hip | 22 +++++++--- examples/bank_conflict/reduce/reduce.hip | 24 +++++++--- .../bank_conflict/synthetic/synthetic.hip | 16 +++++-- .../transpose_scale_add.hip | 24 +++++++--- .../transpose_scale_add_templated.hip | 24 +++++++--- examples/basic/vector_add/vector_add.hip | 44 ++++++++++++------- examples/contention/histogram/histogram.hip | 28 ++++++++---- examples/contention/reduction/reduction.hip | 10 +++++ .../reduction_optimized.hip | 10 +++++ .../simple_reduction/simple_reduction.hip | 30 ++++++++----- src/accordo/_internal/codegen.py | 1 - .../_internal/{hip.py => hip_interop.py} | 0 src/accordo/_internal/ipc/communication.py | 2 +- src/accordo/snapshot.py | 2 +- src/accordo/validator.py | 12 ++--- 17 files changed, 193 insertions(+), 84 deletions(-) rename src/accordo/_internal/{hip.py => hip_interop.py} (100%) diff --git a/examples/access_pattern/uncoalesced/uncoalesced.hip b/examples/access_pattern/uncoalesced/uncoalesced.hip index e7703ae1..82ef614a 100644 --- a/examples/access_pattern/uncoalesced/uncoalesced.hip +++ b/examples/access_pattern/uncoalesced/uncoalesced.hip @@ -53,7 +53,7 @@ __global__ void matrix_transpose(const T* __restrict__ in, } int main() { using T = __hip_bfloat16; - + const int width = 1024; const int height = 1024; const int size = width * height; @@ -97,7 +97,7 @@ int main() { std::cout << (correct ? "Transpose correct ✅" : "Transpose incorrect ❌") << "\n"; - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); return 0; } \ No newline at end of file diff --git a/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip b/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip index 126b698c..2f70de30 100644 --- a/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip +++ b/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip @@ -25,6 +25,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 __global__ void matrixTransposeShared_0(float* out, @@ -106,9 +116,9 @@ void runTranspose(int width, int height) { float* d_in; float* d_out; - hipMalloc(&d_in, width * height * sizeof(float)); - hipMalloc(&d_out, width * height * sizeof(float)); - hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_in, width * height * sizeof(float))); + hip_try(hipMalloc(&d_out, width * height * sizeof(float))); + hip_try(hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice)); dim3 blockSize(TILE_DIM, TILE_DIM); dim3 gridSize((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM); @@ -120,10 +130,10 @@ void runTranspose(int width, int height) { if (status != hipSuccess) { std::terminate(); } - hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost)); - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); } int main() { diff --git a/examples/bank_conflict/matrix_transpose/matrix_transpose.hip b/examples/bank_conflict/matrix_transpose/matrix_transpose.hip index 636fbd08..4d3ebe61 100644 --- a/examples/bank_conflict/matrix_transpose/matrix_transpose.hip +++ b/examples/bank_conflict/matrix_transpose/matrix_transpose.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 __global__ void matrixTransposeShared(float* out, @@ -62,9 +72,9 @@ void runTranspose(int width, int height) { float* d_in; float* d_out; - hipMalloc(&d_in, width * height * sizeof(float)); - hipMalloc(&d_out, width * height * sizeof(float)); - hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_in, width * height * sizeof(float))); + hip_try(hipMalloc(&d_out, width * height * sizeof(float))); + hip_try(hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice)); dim3 blockSize(TILE_DIM, TILE_DIM); dim3 gridSize((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM); @@ -74,10 +84,10 @@ void runTranspose(int width, int height) { if(status != hipSuccess){ std::terminate(); } - hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost)); - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); } int main(int argc, char* argv[]) { diff --git a/examples/bank_conflict/reduce/reduce.hip b/examples/bank_conflict/reduce/reduce.hip index 4cb89c61..483d89fa 100644 --- a/examples/bank_conflict/reduce/reduce.hip +++ b/examples/bank_conflict/reduce/reduce.hip @@ -25,6 +25,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define BLOCK_SIZE 256 __global__ void reduce_kernel(const float *d_in, float *d_out, int n) { @@ -61,22 +71,22 @@ int main() { } float *d_in = nullptr, *d_out = nullptr; - hipMalloc((void **)&d_in, size); - hipMalloc((void **)&d_out, sizeof(float) * (numElements / BLOCK_SIZE)); + hip_try(hipMalloc((void **)&d_in, size)); + hip_try(hipMalloc((void **)&d_out, sizeof(float) * (numElements / BLOCK_SIZE))); - hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice); + hip_try(hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice)); int gridSize = (numElements + BLOCK_SIZE - 1) / BLOCK_SIZE; reduce_kernel<<>>(d_in, d_out, numElements); - hipDeviceSynchronize(); + hip_try(hipDeviceSynchronize()); - hipMemcpy(h_out, d_out, sizeof(float) * gridSize, hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out, d_out, sizeof(float) * gridSize, hipMemcpyDeviceToHost)); printf("First block sum: %f\n", h_out[0]); // Free resources. - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); free(h_in); free(h_out); diff --git a/examples/bank_conflict/synthetic/synthetic.hip b/examples/bank_conflict/synthetic/synthetic.hip index f3af0c91..0b1852b1 100644 --- a/examples/bank_conflict/synthetic/synthetic.hip +++ b/examples/bank_conflict/synthetic/synthetic.hip @@ -24,6 +24,16 @@ SOFTWARE. #include #include + +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + #define BLOCK_SIZE 256 #define LDS_SIZE 256 @@ -52,20 +62,20 @@ int main() { // Allocate memory on the device float* d_out; - hipMalloc(&d_out, BLOCK_SIZE * sizeof(float)); + hip_try(hipMalloc(&d_out, BLOCK_SIZE * sizeof(float))); dim3 blockSize(BLOCK_SIZE); dim3 gridSize(1); hipLaunchKernelGGL(bankConflictKernel, gridSize, blockSize, 0, 0, d_out); // Copy the result back to the host - hipMemcpy(h_out, d_out, BLOCK_SIZE * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out, d_out, BLOCK_SIZE * sizeof(float), hipMemcpyDeviceToHost)); for (int i = 0; i < BLOCK_SIZE; ++i) { std::cout << "h_out[" << i << "] = " << h_out[i] << std::endl; } - hipFree(d_out); + hip_try(hipFree(d_out)); return 0; } diff --git a/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip b/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip index 77f83a90..3926b017 100644 --- a/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip +++ b/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 __global__ void matrixTransposeShared(float* out, @@ -75,9 +85,9 @@ void runKernels(int width, int height) { std::iota(h_data.begin(), h_data.end(), 0.0f); float *d_data, *d_out; - hipMalloc(&d_data, num_elements * sizeof(float)); - hipMalloc(&d_out, num_elements * sizeof(float)); - hipMemcpy(d_data, h_data.data(), num_elements * sizeof(float), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_data, num_elements * sizeof(float))); + hip_try(hipMalloc(&d_out, num_elements * sizeof(float))); + hip_try(hipMemcpy(d_data, h_data.data(), num_elements * sizeof(float), hipMemcpyHostToDevice)); // --- 1. Scale all elements by 2.0 int blockSize = 256; @@ -93,11 +103,11 @@ void runKernels(int width, int height) { matrixTransposeShared<<>>(d_out, d_data, width, height); // --- Copy result back - hipDeviceSynchronize(); - hipMemcpy(h_out.data(), d_out, num_elements * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipDeviceSynchronize()); + hip_try(hipMemcpy(h_out.data(), d_out, num_elements * sizeof(float), hipMemcpyDeviceToHost)); - hipFree(d_data); - hipFree(d_out); + hip_try(hipFree(d_data)); + hip_try(hipFree(d_out)); // Print a few values std::cout << "Result (first 10 elements):\n"; diff --git a/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip b/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip index a0b3aa45..09c6ae0a 100644 --- a/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip +++ b/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 template @@ -76,9 +86,9 @@ void runKernels(int width, int height) { std::iota(h_data.begin(), h_data.end(), static_cast(0)); T *d_data, *d_out; - hipMalloc(&d_data, num_elements * sizeof(T)); - hipMalloc(&d_out, num_elements * sizeof(T)); - hipMemcpy(d_data, h_data.data(), num_elements * sizeof(T), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_data, num_elements * sizeof(T))); + hip_try(hipMalloc(&d_out, num_elements * sizeof(T))); + hip_try(hipMemcpy(d_data, h_data.data(), num_elements * sizeof(T), hipMemcpyHostToDevice)); // --- 1. Scale all elements int blockSize = 256; @@ -93,11 +103,11 @@ void runKernels(int width, int height) { dim3 gridDim((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM); matrixTransposeShared<<>>(d_out, d_data, width, height); - hipDeviceSynchronize(); - hipMemcpy(h_out.data(), d_out, num_elements * sizeof(T), hipMemcpyDeviceToHost); + hip_try(hipDeviceSynchronize()); + hip_try(hipMemcpy(h_out.data(), d_out, num_elements * sizeof(T), hipMemcpyDeviceToHost)); - hipFree(d_data); - hipFree(d_out); + hip_try(hipFree(d_data)); + hip_try(hipFree(d_out)); std::cout << "Result (first 10 elements):\n"; for (int i = 0; i < 10; ++i) diff --git a/examples/basic/vector_add/vector_add.hip b/examples/basic/vector_add/vector_add.hip index ff8da547..a88189f5 100644 --- a/examples/basic/vector_add/vector_add.hip +++ b/examples/basic/vector_add/vector_add.hip @@ -25,6 +25,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + __global__ void vector_add(const float* a, const float* b, float* c, size_t n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) @@ -46,38 +56,38 @@ int main() { } float *d_a, *d_b, *d_c; - hipMalloc(&d_a, size); - hipMalloc(&d_b, size); - hipMalloc(&d_c, size); + hip_try(hipMalloc(&d_a, size)); + hip_try(hipMalloc(&d_b, size)); + hip_try(hipMalloc(&d_c, size)); - hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice); - hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice); + hip_try(hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice)); const int threadsPerBlock = 256; const int blocks = (N + threadsPerBlock - 1) / threadsPerBlock; // Create HIP events hipEvent_t start, stop; - hipEventCreate(&start); - hipEventCreate(&stop); + hip_try(hipEventCreate(&start)); + hip_try(hipEventCreate(&stop)); // Record start - hipEventRecord(start, 0); + hip_try(hipEventRecord(start, 0)); // Launch kernel hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threadsPerBlock), 0, 0, d_a, d_b, d_c, N); // Record stop - hipEventRecord(stop, 0); - hipEventSynchronize(stop); + hip_try(hipEventRecord(stop, 0)); + hip_try(hipEventSynchronize(stop)); // Calculate elapsed time float milliseconds = 0; - hipEventElapsedTime(&milliseconds, start, stop); + hip_try(hipEventElapsedTime(&milliseconds, start, stop)); std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl; // Copy result back to host - hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost)); // Print some results for (size_t i = 0; i < 5; ++i) @@ -87,11 +97,11 @@ int main() { delete[] h_a; delete[] h_b; delete[] h_c; - hipFree(d_a); - hipFree(d_b); - hipFree(d_c); - hipEventDestroy(start); - hipEventDestroy(stop); + hip_try(hipFree(d_a)); + hip_try(hipFree(d_b)); + hip_try(hipFree(d_c)); + hip_try(hipEventDestroy(start)); + hip_try(hipEventDestroy(stop)); return 0; } diff --git a/examples/contention/histogram/histogram.hip b/examples/contention/histogram/histogram.hip index 25bcc996..0ace48f4 100644 --- a/examples/contention/histogram/histogram.hip +++ b/examples/contention/histogram/histogram.hip @@ -5,6 +5,16 @@ #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + using index_type = std::size_t; using real_type = double; @@ -75,17 +85,17 @@ int main(int argc, char* argv[]) { real_type* d_input; real_type* d_output; - hipMalloc(&d_labels, n * sizeof(index_type)); - hipMalloc(&d_flip, n * sizeof(index_type)); - hipMalloc(&d_input, n * sizeof(real_type)); - hipMalloc(&d_output, k * sizeof(real_type)); + hip_try(hipMalloc(&d_labels, n * sizeof(index_type))); + hip_try(hipMalloc(&d_flip, n * sizeof(index_type))); + hip_try(hipMalloc(&d_input, n * sizeof(real_type))); + hip_try(hipMalloc(&d_output, k * sizeof(real_type))); - hipMemcpy(d_labels, h_labels.data(), n * sizeof(index_type), hipMemcpyHostToDevice); - hipMemcpy(d_flip, h_flip.data(), n * sizeof(index_type), hipMemcpyHostToDevice); - hipMemcpy(d_input, h_input.data(), n * sizeof(real_type), hipMemcpyHostToDevice); - hipMemcpy(d_output, h_output.data(), k * sizeof(real_type), hipMemcpyHostToDevice); + hip_try(hipMemcpy(d_labels, h_labels.data(), n * sizeof(index_type), hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_flip, h_flip.data(), n * sizeof(index_type), hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_input, h_input.data(), n * sizeof(real_type), hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_output, h_output.data(), k * sizeof(real_type), hipMemcpyHostToDevice)); histogram(n, k, d_labels, d_flip, d_input, d_output); - hipDeviceSynchronize(); + hip_try(hipDeviceSynchronize()); } \ No newline at end of file diff --git a/examples/contention/reduction/reduction.hip b/examples/contention/reduction/reduction.hip index c6d2948e..78886216 100644 --- a/examples/contention/reduction/reduction.hip +++ b/examples/contention/reduction/reduction.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + __global__ void reduction_kernel(const float* input, float* result, std::size_t count) { const auto thread_id = threadIdx.x + blockIdx.x * blockDim.x; if (thread_id < count) { diff --git a/examples/contention/reduction_optimized/reduction_optimized.hip b/examples/contention/reduction_optimized/reduction_optimized.hip index 4772bbdf..72db0f73 100644 --- a/examples/contention/reduction_optimized/reduction_optimized.hip +++ b/examples/contention/reduction_optimized/reduction_optimized.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + __global__ void reduction_kernel(const float* input, float* result, std::size_t count) { extern __shared__ float shared_data[]; const auto thread_id = threadIdx.x + blockIdx.x * blockDim.x; diff --git a/examples/contention/simple_reduction/simple_reduction.hip b/examples/contention/simple_reduction/simple_reduction.hip index 1339f9a0..92f9535b 100644 --- a/examples/contention/simple_reduction/simple_reduction.hip +++ b/examples/contention/simple_reduction/simple_reduction.hip @@ -28,6 +28,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + using data_t = double; __global__ void reduction_kernel(const data_t* input, data_t* result, std::size_t count) { @@ -45,33 +55,33 @@ int main() { data_t* d_input = nullptr; data_t* d_result = nullptr; - hipMalloc(&d_input, count * sizeof(data_t)); - hipMalloc(&d_result, sizeof(data_t)); + hip_try(hipMalloc(&d_input, count * sizeof(data_t))); + hip_try(hipMalloc(&d_result, sizeof(data_t))); std::vector h_input(count, 1); - hipMemcpy(d_input, h_input.data(), count * sizeof(data_t), hipMemcpyHostToDevice); - hipMemset(d_result, 0, sizeof(data_t)); + hip_try(hipMemcpy(d_input, h_input.data(), count * sizeof(data_t), hipMemcpyHostToDevice)); + hip_try(hipMemset(d_result, 0, sizeof(data_t))); std::cout << "input: " << d_input << std::endl; std::cout << "result: " << d_result << std::endl; reduction_kernel<<>>(d_input, d_result, count); - hipDeviceSynchronize(); + hip_try(hipDeviceSynchronize()); data_t h_result = 0; - hipMemcpy(&h_result, d_result, sizeof(data_t), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(&h_result, d_result, sizeof(data_t), hipMemcpyDeviceToHost)); if (h_result != count) { std::cout << "Kernel failed. Expected: " << count << ", Got: " << h_result << "\n"; - hipFree(d_input); - hipFree(d_result); + hip_try(hipFree(d_input)); + hip_try(hipFree(d_result)); return -1; } else { std::cout << "Success!"; } std::cout << std::endl; - hipFree(d_input); - hipFree(d_result); + hip_try(hipFree(d_input)); + hip_try(hipFree(d_result)); return 0; } diff --git a/src/accordo/_internal/codegen.py b/src/accordo/_internal/codegen.py index abbae9f9..3fc9834a 100644 --- a/src/accordo/_internal/codegen.py +++ b/src/accordo/_internal/codegen.py @@ -1,7 +1,6 @@ """Code generation for Accordo C++ header files.""" import logging -from pathlib import Path def generate_kernel_header(args: list[str], additional_includes: list[str] = None) -> str: diff --git a/src/accordo/_internal/hip.py b/src/accordo/_internal/hip_interop.py similarity index 100% rename from src/accordo/_internal/hip.py rename to src/accordo/_internal/hip_interop.py diff --git a/src/accordo/_internal/ipc/communication.py b/src/accordo/_internal/ipc/communication.py index 200c8a2a..f9df057c 100644 --- a/src/accordo/_internal/ipc/communication.py +++ b/src/accordo/_internal/ipc/communication.py @@ -9,7 +9,7 @@ import ml_dtypes import numpy as np -from ..hip import memcpy_d2h, open_ipc_handle +from ..hip_interop import memcpy_d2h, open_ipc_handle def read_ipc_handles(args, ipc_file_name): diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py index e496dba0..42b34c6d 100644 --- a/src/accordo/snapshot.py +++ b/src/accordo/snapshot.py @@ -44,7 +44,7 @@ def summary(self) -> str: """Get a detailed summary of the snapshot.""" binary_str = " ".join(self.binary) lines = [ - f"Snapshot Summary:", + "Snapshot Summary:", f" Binary: {binary_str}", f" Working Directory: {self.working_directory}", f" Execution Time: {self.execution_time_ms:.2f}ms", diff --git a/src/accordo/validator.py b/src/accordo/validator.py index 94827948..d88fae86 100644 --- a/src/accordo/validator.py +++ b/src/accordo/validator.py @@ -10,12 +10,12 @@ import numpy as np +from ._internal.codegen import generate_kernel_header +from ._internal.ipc.communication import get_kern_arg_data, send_response from .config import ValidationConfig -from .exceptions import AccordoBuildError, AccordoProcessError, AccordoTimeoutError, AccordoValidationError +from .exceptions import AccordoBuildError, AccordoProcessError, AccordoTimeoutError from .result import ArrayMismatch, ValidationResult from .snapshot import Snapshot -from ._internal.codegen import generate_kernel_header -from ._internal.ipc.communication import get_kern_arg_data, send_response class _TimeoutException(Exception): @@ -361,12 +361,12 @@ def _run_instrumented_app( process_pid=process_pid, baseline_time_ms=baseline_time_ms, ) - except TimeoutError as e: + except TimeoutError: # Kill the process if it timed out try: os.kill(process_pid, 9) - except: - pass + except (OSError, ProcessLookupError): + pass # Process already dead raise # Send completion response From 9ab60ee27cbfe7b392a8f22707678b1ad81d65db Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" Date: Mon, 3 Nov 2025 08:05:49 +0000 Subject: [PATCH 06/15] Apply Ruff auto-fixes --- src/accordo/_internal/__init__.py | 1 - src/accordo/_internal/codegen.py | 1 - src/accordo/_internal/hip_interop.py | 1 - src/accordo/_internal/ipc/__init__.py | 1 - src/accordo/_internal/ipc/communication.py | 1 - src/accordo/config.py | 1 - src/accordo/exceptions.py | 1 - src/accordo/result.py | 1 - src/accordo/snapshot.py | 1 - src/accordo/validator.py | 12 ++++-------- src/intelliperf/formulas/formula_base.py | 14 +++----------- 11 files changed, 7 insertions(+), 28 deletions(-) diff --git a/src/accordo/_internal/__init__.py b/src/accordo/_internal/__init__.py index edfbb039..7a0b6b33 100644 --- a/src/accordo/_internal/__init__.py +++ b/src/accordo/_internal/__init__.py @@ -1,2 +1 @@ """Internal implementation details for Accordo. Not part of public API.""" - diff --git a/src/accordo/_internal/codegen.py b/src/accordo/_internal/codegen.py index 3fc9834a..fed8524a 100644 --- a/src/accordo/_internal/codegen.py +++ b/src/accordo/_internal/codegen.py @@ -48,4 +48,3 @@ def generate_kernel_header(args: list[str], additional_includes: list[str] = Non logging.debug(f"Generated header file: {header_path}") logging.debug(f"Header content: {header_content}") return header_path - diff --git a/src/accordo/_internal/hip_interop.py b/src/accordo/_internal/hip_interop.py index b05bfc0d..d30c822f 100644 --- a/src/accordo/_internal/hip_interop.py +++ b/src/accordo/_internal/hip_interop.py @@ -104,4 +104,3 @@ def memcpy_d2h(ptr, num_elements_to_copy, dtype): ) ) return host_array - diff --git a/src/accordo/_internal/ipc/__init__.py b/src/accordo/_internal/ipc/__init__.py index 9af22b27..c8ef421d 100644 --- a/src/accordo/_internal/ipc/__init__.py +++ b/src/accordo/_internal/ipc/__init__.py @@ -1,2 +1 @@ """IPC communication modules for Accordo.""" - diff --git a/src/accordo/_internal/ipc/communication.py b/src/accordo/_internal/ipc/communication.py index f9df057c..251f53f8 100644 --- a/src/accordo/_internal/ipc/communication.py +++ b/src/accordo/_internal/ipc/communication.py @@ -163,4 +163,3 @@ def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30, pr results.append(host_array) return results - diff --git a/src/accordo/config.py b/src/accordo/config.py index 6bcb5278..43efc69f 100644 --- a/src/accordo/config.py +++ b/src/accordo/config.py @@ -104,4 +104,3 @@ def get_arg_types(self) -> list[str]: def get_arg_names(self) -> list[str]: """Get list of argument names.""" return [arg.name for arg in self.kernel_args] - diff --git a/src/accordo/exceptions.py b/src/accordo/exceptions.py index f0f739e6..b9476fd4 100644 --- a/src/accordo/exceptions.py +++ b/src/accordo/exceptions.py @@ -33,4 +33,3 @@ class AccordoValidationError(AccordoError): """Raised when array validation fails.""" pass - diff --git a/src/accordo/result.py b/src/accordo/result.py index 5664acda..9a824b66 100644 --- a/src/accordo/result.py +++ b/src/accordo/result.py @@ -98,4 +98,3 @@ def summary(self) -> str: def __str__(self) -> str: """String representation.""" return self.summary() - diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py index 42b34c6d..6b08a506 100644 --- a/src/accordo/snapshot.py +++ b/src/accordo/snapshot.py @@ -55,4 +55,3 @@ def summary(self) -> str: lines.append(f" Array {i}: shape={arr.shape}, dtype={arr.dtype}") return "\n".join(lines) - diff --git a/src/accordo/validator.py b/src/accordo/validator.py index d88fae86..13f67723 100644 --- a/src/accordo/validator.py +++ b/src/accordo/validator.py @@ -20,6 +20,7 @@ class _TimeoutException(Exception): """Internal exception for timeout handling.""" + pass @@ -196,10 +197,7 @@ def capture_snapshot( try: start_time = time.time() result_arrays = self._run_instrumented_app( - binary, - working_directory, - label="snapshot", - baseline_time_ms=None + binary, working_directory, label="snapshot", baseline_time_ms=None ) signal.alarm(0) # Cancel alarm on success execution_time_ms = (time.time() - start_time) * 1000 @@ -214,9 +212,8 @@ def capture_snapshot( signal.alarm(0) logging.error(f"Snapshot capture timed out after {timeout_seconds}s") raise AccordoTimeoutError( - f"Snapshot capture timed out after {timeout_seconds}s. " - "This may indicate a GPU crash or hung process.", - timeout_seconds=timeout_seconds + f"Snapshot capture timed out after {timeout_seconds}s. This may indicate a GPU crash or hung process.", + timeout_seconds=timeout_seconds, ) except TimeoutError as e: signal.alarm(0) @@ -451,4 +448,3 @@ def _validate_results( matched_arrays=matched_arrays, execution_time_ms=execution_times, ) - diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 88d1b190..2c5d1b35 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -41,7 +41,6 @@ from intelliperf import __version__ from intelliperf.core.application import Application from intelliperf.core.logger import Logger -from intelliperf.utils.env import get_accordo_path from intelliperf.utils.process import capture_subprocess_output, exit_on_fail @@ -461,9 +460,7 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole logging.debug("Capturing reference snapshot (will be cached)") self._reference_snapshot = self._accordo_validator.capture_snapshot( - binary=reference_binary, - working_directory=working_dir, - timeout_seconds=30 + binary=reference_binary, working_directory=working_dir, timeout_seconds=30 ) logging.debug(f"Reference snapshot captured in {self._reference_snapshot.execution_time_ms:.2f}ms") except Exception as e: @@ -484,17 +481,12 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole logging.debug("Capturing optimized snapshot") optimized_snapshot = self._accordo_validator.capture_snapshot( - binary=optimized_binary, - working_directory=working_dir, - timeout_seconds=opt_timeout + binary=optimized_binary, working_directory=working_dir, timeout_seconds=opt_timeout ) logging.debug(f"Optimized snapshot captured in {optimized_snapshot.execution_time_ms:.2f}ms") # Compare snapshots - validation_result = self._accordo_validator.compare_snapshots( - self._reference_snapshot, - optimized_snapshot - ) + validation_result = self._accordo_validator.compare_snapshots(self._reference_snapshot, optimized_snapshot) if validation_result.is_valid: logging.debug("Validation succeeded.") From 0ff197a096a7fec4abcaef201b94454208221e49 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 02:19:11 -0600 Subject: [PATCH 07/15] Fix maximize --- src/intelliperf/formulas/atomic_contention.py | 2 +- src/intelliperf/formulas/bank_conflict.py | 4 ++-- src/intelliperf/formulas/formula_base.py | 12 ++++++++++-- 3 files changed, 13 insertions(+), 5 deletions(-) diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index 0ec0b732..5105114a 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -263,7 +263,7 @@ def __init__( self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="latency_improvement", - maximize=True, + maximize=False, # False = minimize raw metric (latency) before_metric="unoptimized_lat", after_metric="optimized_lat", ) diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index b096d768..e01d8ca5 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -270,12 +270,12 @@ def __init__( self.success = False # Initialize optimization tracker - # Bank conflict optimization maximizes conflict reduction (minimize bank conflicts) + # Bank conflict optimization minimizes conflicts (lower is better) # Automatically calculates conflict_improvement from unoptimized_conflicts / optimized_conflicts self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="conflict_improvement", - maximize=True, + maximize=False, # False = minimize raw metric (conflicts) before_metric="unoptimized_conflicts", after_metric="optimized_conflicts", ) diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 2c5d1b35..4a40f0c9 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -99,9 +99,17 @@ def add_step( if self.before_metric and self.after_metric: before = metrics.get(self.before_metric, 0) after = metrics.get(self.after_metric, 0) - if before != 0: - improvement = after / before if after != 0 else 1.0 + if before != 0 and after != 0: + # Calculate improvement ratio based on whether we're maximizing or minimizing the raw metric + # For conflicts/latency (maximize=False, lower is better): improvement = before / after + # Example: 3.5 conflicts → 1.0 conflicts = 3.5 / 1.0 = 3.5x improvement + # For coalescing (maximize=True, higher is better): improvement = after / before + # Example: 50% → 75% = 75 / 50 = 1.5x improvement + improvement = before / after if not self.maximize else after / before metrics[self.primary_metric] = improvement + elif before != 0: + # If after is 0, set improvement to 1.0 (no change) + metrics[self.primary_metric] = 1.0 step = OptimizationStep( iteration=self.current_iteration, From 58c44589eca5282558540ec4e4ea9728e4e32ed0 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 02:21:15 -0600 Subject: [PATCH 08/15] Fix conflict_improvement calculation and JSON field ordering --- src/intelliperf/formulas/formula_base.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 4a40f0c9..44c58a80 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -637,9 +637,9 @@ def write_results( "version": __version__, "optimized": self._optimization_results, "initial": self._initial_profiler_results, + **additional_results, "report_message": self.optimization_report, "bottleneck_report": self.bottleneck_report, - **additional_results, "diff": self.compute_diff(self.current_kernel_files), } if self.in_place: From 7509f7f72c104a797e0aa399911577c39574f256 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 03:38:20 -0600 Subject: [PATCH 09/15] Log source --- src/intelliperf/__main__.py | 4 +++ src/intelliperf/core/logger.py | 26 +++++++++++++++++++ src/intelliperf/formulas/atomic_contention.py | 7 ++--- src/intelliperf/formulas/bank_conflict.py | 8 +++--- src/intelliperf/formulas/formula_base.py | 24 +++++++++++++++++ src/intelliperf/formulas/memory_access.py | 8 +++--- 6 files changed, 66 insertions(+), 11 deletions(-) diff --git a/src/intelliperf/__main__.py b/src/intelliperf/__main__.py index bf2d0926..f66c1c9f 100644 --- a/src/intelliperf/__main__.py +++ b/src/intelliperf/__main__.py @@ -258,6 +258,10 @@ def main(): optimizer = formula(**optimizer_args) + # Store trace path in logger for use in iteration logging + if hasattr(optimizer, 'get_logger'): + optimizer.get_logger().trace_path = args.trace_path + # Helper function to flush logs if tracing is enabled def flush_logs_if_enabled(): if hasattr(optimizer, "get_logger") and args.trace_path: diff --git a/src/intelliperf/core/logger.py b/src/intelliperf/core/logger.py index d52f9170..226a850d 100644 --- a/src/intelliperf/core/logger.py +++ b/src/intelliperf/core/logger.py @@ -190,6 +190,32 @@ def get_run_summary(self) -> Dict[str, Any]: "events": self.buffer, } + def save_iteration_code(self, kernel_name: str, iteration_num: int, code_content: str) -> str: + """ + Save iteration code to a file in the trace directory. + + Args: + kernel_name: Name of the kernel being optimized + iteration_num: Iteration number + code_content: The code content to save + + Returns: + The path to the saved file + """ + # Get output directory from trace_path if available + if hasattr(self, 'trace_path') and self.trace_path: + output_dir = os.path.abspath(self.trace_path) + os.makedirs(output_dir, exist_ok=True) + else: + output_dir = os.path.abspath(".") + + iteration_file = os.path.join(output_dir, f"{kernel_name}_iteration_{iteration_num}.hip") + with open(iteration_file, "w") as f: + f.write(code_content) + logging.info(f"Saved iteration {iteration_num} code to {iteration_file}") + + return iteration_file + def flush(self, output_file: Optional[str] = None) -> bool: """ Flush logs to output targets with error handling. diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index 5105114a..b562446c 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -541,9 +541,10 @@ def optimize_pass( }, ) - with open(kernel_file, "w") as f: - f.write(optimized_file_content) + # Write and log optimized code immediately after LLM generation + self.write_and_log_optimized_code(kernel_file, optimized_file_content) logging.debug(f"Optimized file content: {optimized_file_content}") + return Result( success=True, asset={ @@ -786,9 +787,9 @@ def write_results(self, output_file: str = None): super().write_results( output_file=output_file, additional_results={ + "optimization_history": self.optimization_tracker.to_dict(), "formula": "atomicContention", "success": self.success, - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index e01d8ca5..f72436c2 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -629,9 +629,9 @@ def optimize_pass( }, ) - with open(kernel_file, "w") as f: - f.write(optimized_file_content) - logging.debug(f"Optimized file content: {optimized_file_content}") + # Write and log optimized code immediately after LLM generation + self.write_and_log_optimized_code(kernel_file, optimized_file_content) + return Result( success=True, asset={ @@ -863,9 +863,9 @@ def write_results(self, output_file: str = None): super().write_results( output_file=output_file, additional_results={ + "optimization_history": self.optimization_tracker.to_dict(), "formula": "bankConflict", "success": self.success, - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 44c58a80..1648f6f7 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -574,6 +574,30 @@ def postprocess_llm_code(self, optimized_file_content: str) -> str: return content + def write_and_log_optimized_code(self, kernel_file: str, optimized_code: str) -> None: + """ + Write optimized code to file and log it for future reference. + + This should be called immediately after LLM generates code, regardless of whether + it will compile or pass validation. + + Args: + kernel_file: Path to the kernel file to write + optimized_code: The optimized code content + """ + # Write the code to the kernel file + with open(kernel_file, "w") as f: + f.write(optimized_code) + + # Automatically detect iteration number from optimization tracker + iteration_num = len(self.optimization_tracker.steps) if hasattr(self, 'optimization_tracker') else 0 + + # Log the iteration code immediately + kernel_name = get_kernel_name(self.current_kernel_signature if hasattr(self, 'current_kernel_signature') else "kernel") + self.get_logger().save_iteration_code(kernel_name, iteration_num, optimized_code) + + logging.debug(f"Wrote and logged optimized code to {kernel_file} (iteration {iteration_num})") + def compute_diff(self, filepaths: list[str]) -> str: diffs = [] for filepath in filepaths: diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index 3d5c2000..a1485b79 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -540,9 +540,9 @@ def optimize_pass( }, ) - with open(kernel_file, "w") as f: - f.write(optimized_file_content) - logging.debug(f"Optimized file content: {optimized_file_content}") + # Write and log optimized code immediately after LLM generation + self.write_and_log_optimized_code(kernel_file, optimized_file_content) + return Result( success=True, asset={ @@ -771,9 +771,9 @@ def write_results(self, output_file: str = None): super().write_results( output_file=output_file, additional_results={ + "optimization_history": self.optimization_tracker.to_dict(), "formula": "memoryAccess", "success": self.success, - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) From 61fe7660cde6baebd8ea45bea2deae89b13092c2 Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" Date: Mon, 3 Nov 2025 09:38:51 +0000 Subject: [PATCH 10/15] Apply Ruff auto-fixes --- src/intelliperf/__main__.py | 2 +- src/intelliperf/core/logger.py | 2 +- src/intelliperf/formulas/formula_base.py | 6 ++++-- 3 files changed, 6 insertions(+), 4 deletions(-) diff --git a/src/intelliperf/__main__.py b/src/intelliperf/__main__.py index f66c1c9f..d9e5e9f0 100644 --- a/src/intelliperf/__main__.py +++ b/src/intelliperf/__main__.py @@ -259,7 +259,7 @@ def main(): optimizer = formula(**optimizer_args) # Store trace path in logger for use in iteration logging - if hasattr(optimizer, 'get_logger'): + if hasattr(optimizer, "get_logger"): optimizer.get_logger().trace_path = args.trace_path # Helper function to flush logs if tracing is enabled diff --git a/src/intelliperf/core/logger.py b/src/intelliperf/core/logger.py index 226a850d..65da6ed8 100644 --- a/src/intelliperf/core/logger.py +++ b/src/intelliperf/core/logger.py @@ -203,7 +203,7 @@ def save_iteration_code(self, kernel_name: str, iteration_num: int, code_content The path to the saved file """ # Get output directory from trace_path if available - if hasattr(self, 'trace_path') and self.trace_path: + if hasattr(self, "trace_path") and self.trace_path: output_dir = os.path.abspath(self.trace_path) os.makedirs(output_dir, exist_ok=True) else: diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 1648f6f7..4066b95a 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -590,10 +590,12 @@ def write_and_log_optimized_code(self, kernel_file: str, optimized_code: str) -> f.write(optimized_code) # Automatically detect iteration number from optimization tracker - iteration_num = len(self.optimization_tracker.steps) if hasattr(self, 'optimization_tracker') else 0 + iteration_num = len(self.optimization_tracker.steps) if hasattr(self, "optimization_tracker") else 0 # Log the iteration code immediately - kernel_name = get_kernel_name(self.current_kernel_signature if hasattr(self, 'current_kernel_signature') else "kernel") + kernel_name = get_kernel_name( + self.current_kernel_signature if hasattr(self, "current_kernel_signature") else "kernel" + ) self.get_logger().save_iteration_code(kernel_name, iteration_num, optimized_code) logging.debug(f"Wrote and logged optimized code to {kernel_file} (iteration {iteration_num})") From 2994f9687b7f7e063f2e04d9041d5b3f2f5d2084 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Sun, 2 Nov 2025 23:00:29 -0800 Subject: [PATCH 11/15] Unify counters report (#155) Co-authored-by: github-actions[bot] --- .gitignore | 2 + pyproject.toml | 4 +- src/accordo/python/communicate.py | 140 +++++++++++++++++++++----- src/intelliperf/formulas/swizzling.py | 27 ++--- 4 files changed, 124 insertions(+), 49 deletions(-) diff --git a/.gitignore b/.gitignore index 41639bab..5132677b 100644 --- a/.gitignore +++ b/.gitignore @@ -14,3 +14,5 @@ external/ intelliperf_env/ trace/ .build/ + +.rocprofv3/ \ No newline at end of file diff --git a/pyproject.toml b/pyproject.toml index 3f08b319..99bc8197 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -12,9 +12,9 @@ authors = [ license = { text = "MIT" } readme = "README.md" -requires-python = ">=3.10" +requires-python = ">=3.9" -# Python dependencies +# Python dependencies dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==2.6.27", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]", "rpds-py"] [tool.setuptools] diff --git a/src/accordo/python/communicate.py b/src/accordo/python/communicate.py index ba433b99..d24760ec 100644 --- a/src/accordo/python/communicate.py +++ b/src/accordo/python/communicate.py @@ -23,10 +23,13 @@ ################################################################################ import ctypes +import errno import logging +import math import os import stat import sys +import threading import time import ml_dtypes @@ -36,6 +39,34 @@ from hip import memcpy_d2h, open_ipc_handle +def run_with_timeout(func, timeout_seconds, *args, **kwargs): + """Cross-platform timeout wrapper using threading. + + Runs func in a thread and raises TimeoutError if it doesn't complete in time. + """ + result = [None] + exception = [None] + + def target(): + try: + result[0] = func(*args, **kwargs) + except Exception as e: + exception[0] = e + + thread = threading.Thread(target=target, daemon=True) + thread.start() + thread.join(timeout=timeout_seconds) + + if thread.is_alive(): + # Thread is still running - timeout occurred + raise TimeoutError(f"Operation timed out after {timeout_seconds} seconds") + + if exception[0] is not None: + raise exception[0] + + return result[0] + + def read_ipc_handles(args, ipc_file_name): count = sum(1 for arg in args if "*" in arg and "const" not in arg) @@ -45,7 +76,6 @@ def read_ipc_handles(args, ipc_file_name): while len(handles) < count: if not os.path.exists(ipc_file_name): - logging.debug("Waiting for IPC file...") time.sleep(0.1) continue @@ -71,18 +101,18 @@ def read_ipc_handles(args, ipc_file_name): size_value = int.from_bytes(size_data, byteorder="little") sizes.append(size_value) - logging.debug("Final IPC Handle (hex):") - for i in range(0, len(handle_np), 16): - chunk = handle_np[i : i + 16] - logging.debug(" ".join(f"{b:02x}" for b in chunk)) - - logging.debug(f"Corresponding Pointer Size: {size_value} bytes") + # Verbose IPC handle debugging (only when new handle received) + if logging.getLogger().isEnabledFor(logging.DEBUG): + logging.debug("Final IPC Handle (hex):") + for i in range(0, len(handle_np), 16): + chunk = handle_np[i : i + 16] + logging.debug(" ".join(f"{b:02x}" for b in chunk)) + logging.debug(f"Corresponding Pointer Size: {size_value} bytes") if len(handles) < count: - logging.debug(f"Waiting for {count - len(handles)} more IPC handles...") + # Don't spam logs in hot loop - removed logging.debug here time.sleep(0.1) - # logging.debug(f"Successfully read {len(handles)} IPC handles and sizes.") return handles, sizes @@ -91,27 +121,83 @@ def send_response(pipe_name): fifo.write("done\n") -def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30): +def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30, process_pid=None, baseline_time_ms=None): logging.debug(f"pipe_name: {pipe_name}") logging.debug(f"get_kern_arg_data args: {args}") logging.debug(f"ipc_file_name: {ipc_file_name}") - if not os.path.exists(pipe_name): - os.mkfifo(pipe_name) - os.chmod(pipe_name, stat.S_IRWXU | stat.S_IRWXG | stat.S_IRWXO) - - start_time = time.time() - with open(pipe_name, "rb") as fifo: # noqa: F841 - while True: - if time.time() - start_time > ipc_timeout_seconds: - raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data") - - try: - ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) - break - except Exception as e: - if time.time() - start_time > ipc_timeout_seconds: - raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data: {str(e)}") - time.sleep(0.1) + + # Calculate dynamic timeout based on baseline performance + if baseline_time_ms is not None and baseline_time_ms > 0: + # 2x baseline, rounded up to next second, minimum 3 seconds + ipc_timeout_seconds = max(3, math.ceil(baseline_time_ms / 1000.0 * 2.0)) + logging.debug(f"Using dynamic timeout: {ipc_timeout_seconds}s (2x baseline of {baseline_time_ms}ms)") + else: + logging.debug(f"Using default timeout: {ipc_timeout_seconds}s (no baseline available)") + + def _do_ipc_work(): + """Inner function that does the actual IPC work - wrapped with timeout""" + fifo = None + try: + if not os.path.exists(pipe_name): + os.mkfifo(pipe_name) + os.chmod(pipe_name, stat.S_IRWXU | stat.S_IRWXG | stat.S_IRWXO) + + # Try to open the pipe, checking if the process is still alive + while fifo is None: + # Check if the process is still alive (crash detection, not timeout) + if process_pid is not None: + try: + os.kill(process_pid, 0) # Signal 0 just checks if process exists + except OSError: + raise RuntimeError( + f"Accordo process (PID {process_pid}) crashed or terminated before opening pipe. Check for segfaults or GPU memory access errors." + ) + + try: + # Try non-blocking open + fd = os.open(pipe_name, os.O_RDONLY | os.O_NONBLOCK) + fifo = os.fdopen(fd, "rb") + break + except OSError as e: + if e.errno == errno.ENXIO: # ENXIO - no writer connected yet + time.sleep(0.1) + continue + else: + raise + + # Read IPC handles + while True: + # Check if the process is still alive (crash detection, not timeout) + if process_pid is not None: + try: + os.kill(process_pid, 0) + except OSError: + raise RuntimeError( + f"Accordo process (PID {process_pid}) crashed or terminated during execution. Check for segfaults or GPU memory access errors." + ) + + try: + ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) + break + except Exception: + # For non-timeout exceptions, retry with a short sleep + time.sleep(0.1) + + return ipc_handles, ptr_sizes + + finally: + if fifo: + fifo.close() + + # Run IPC work with timeout wrapper (cross-platform) + try: + ipc_handles, ptr_sizes = run_with_timeout(_do_ipc_work, ipc_timeout_seconds) + except TimeoutError: + # Enhance timeout message with context + timeout_msg = f"Timeout after {ipc_timeout_seconds} seconds during IPC communication" + if baseline_time_ms is not None: + timeout_msg += f" (baseline: {baseline_time_ms}ms, 2x timeout: {ipc_timeout_seconds}s). Code may be correct but too slow to be worth profiling." + raise TimeoutError(timeout_msg) type_map = { "double*": ctypes.c_double, diff --git a/src/intelliperf/formulas/swizzling.py b/src/intelliperf/formulas/swizzling.py index 93d6987a..cf5f5430 100644 --- a/src/intelliperf/formulas/swizzling.py +++ b/src/intelliperf/formulas/swizzling.py @@ -27,7 +27,6 @@ import logging import os import stat -import sys import dspy @@ -37,7 +36,6 @@ Formula_Base, Result, filter_json_field, - get_kernel_name, ) from intelliperf.utils.env import get_llm_api_key @@ -256,22 +254,8 @@ def optimize_pass( kernel = filtered_report_card[0]["kernel"] files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - - logging.debug(f"Kernel name: {kernel_name}") - kernel_file = None - for file in files: - if os.path.exists(file): - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - logging.error(f"Kernel file not found for kernel {kernel}") - sys.exit(1) - else: - logging.debug(f"Kernel file found for kernel {kernel}: {kernel_file}") + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) + logging.debug(f"Kernel file found for kernel {kernel}: {kernel_file}") # Stage 1: Memory access pattern analysis (only run once) if not self.memory_analysis_done: @@ -287,7 +271,7 @@ def optimize_pass( self.bottleneck_report = ( f"L2 Cache Locality Detection: IntelliPerf identified suboptimal L2 cache hit rate " - f"in kernel `{kernel_name}`. Poor cache locality occurs when " + f"in kernel `{kernel}`. Poor cache locality occurs when " f"blocks accessing related memory are scheduled to different XCDs with separate L2 caches, " f"reducing overall cache effectiveness." ) @@ -532,7 +516,10 @@ def performance_validation_pass(self) -> Result: if self.current_iteration < self.max_iterations: self.current_summary = self.optimization_report # Always return success=False to continue iterating - return Result(success=False, error_report=self.best_iteration_report) + error_msg = ( + self.best_iteration_report if self.best_iteration_report else "Continuing optimization iterations..." + ) + return Result(success=False, error_report=error_msg) return Result(success=True, asset={"log": self.best_iteration_report}) From f2c3160bbf1ece2c01d5f19b5cc506ac171fe478 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 04:00:57 -0600 Subject: [PATCH 12/15] Add find_kernel_file method from main to all formulas while keeping Accordo improvements --- src/intelliperf/formulas/atomic_contention.py | 16 +---- src/intelliperf/formulas/bank_conflict.py | 17 +---- src/intelliperf/formulas/formula_base.py | 64 +++++++++++++++++++ src/intelliperf/formulas/memory_access.py | 20 +----- 4 files changed, 67 insertions(+), 50 deletions(-) diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index b562446c..f0b1518a 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -434,21 +434,7 @@ def optimize_pass( self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - kernel_file = None - unoptimized_file_content = None - for file in files: - project_dir = os.path.abspath(self._application.get_project_directory()) - file_path = os.path.abspath(file) - isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir - if os.path.exists(file) and isfile_in_project: - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - return Result(success=False, error_report="Kernel file not found.") + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) # Build problem description problem_description = ( diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index f72436c2..a4949fad 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -522,22 +522,7 @@ def optimize_pass( self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - kernel_file = None - - unoptimized_file_content = None - for file in files: - project_dir = os.path.abspath(self._application.get_project_directory()) - file_path = os.path.abspath(file) - isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir - if os.path.exists(file) and isfile_in_project: - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - return Result(success=False, error_report="Kernel file not found.") + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) # Build problem description problem_description = ( diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 4066b95a..6c97a495 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -600,6 +600,70 @@ def write_and_log_optimized_code(self, kernel_file: str, optimized_code: str) -> logging.debug(f"Wrote and logged optimized code to {kernel_file} (iteration {iteration_num})") + def find_kernel_file(self, files: list, kernel: str) -> tuple: + """ + Find the kernel file containing the given kernel name from a list of files. + + Args: + files: List of file paths to search + kernel: Kernel signature to find + + Returns: + tuple: (kernel_file_path, file_content) or (None, None) if not found + + Note: + - Validates files exist and are within the project directory + - Logs warnings for invalid files + - Exits with sys.exit(1) if kernel file not found after checking all files + """ + kernel_name = get_kernel_name(kernel) + logging.debug(f"Searching for kernel: {kernel_name}") + + kernel_file = None + unoptimized_file_content = None + project_dir = os.path.abspath(self._application.get_project_directory()) + + for file in files: + file_path = os.path.abspath(file) + + # Check if file exists + if not os.path.exists(file): + logging.warning(f"File {file} does not exist") + continue + + # Check if file is in project directory + try: + isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir + except ValueError: + # Happens when paths are on different drives (Windows) + isfile_in_project = False + + if not isfile_in_project: + logging.warning(f"File {file} is not in the project") + continue + + # Try to read file and find kernel + try: + with open(file, "r") as f: + unoptimized_file_content = f.read() + if kernel_name in unoptimized_file_content: + kernel_file = file + break + except Exception as e: + logging.error(f"Error reading file {file}: {e}") + continue + + # If kernel file not found, log error and exit + if kernel_file is None: + logging.error(f"Kernel file not found for kernel {kernel}") + logging.error(f"Kernel name: {kernel_name}") + logging.error(f"Files searched: {files}") + if unoptimized_file_content: + logging.error(f"Last file content (first 200 chars): {unoptimized_file_content[:200]}") + sys.exit(1) + + return kernel_file, unoptimized_file_content + def compute_diff(self, filepaths: list[str]) -> str: diffs = [] for filepath in filepaths: diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index a1485b79..c35e6f39 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -430,25 +430,7 @@ def optimize_pass( self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - - logging.debug(f"Kernel name: {kernel_name}") - kernel_file = None - unoptimized_file_content = None - for file in files: - project_dir = os.path.abspath(self._application.get_project_directory()) - file_path = os.path.abspath(file) - isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir - - if os.path.exists(file) and isfile_in_project: - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - logging.error(f"Kernel file not found for kernel {kernel}") - sys.exit(1) + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) # Build problem description problem_description = ( From 0979991cbc2ff5467f77bb86b41ab2aa0be789bc Mon Sep 17 00:00:00 2001 From: "github-actions[bot]" Date: Mon, 3 Nov 2025 10:01:22 +0000 Subject: [PATCH 13/15] Apply Ruff auto-fixes --- src/intelliperf/formulas/atomic_contention.py | 2 -- src/intelliperf/formulas/bank_conflict.py | 1 - src/intelliperf/formulas/memory_access.py | 2 -- 3 files changed, 5 deletions(-) diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index f0b1518a..9cfbb1f6 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -24,7 +24,6 @@ import json import logging -import os import dspy @@ -33,7 +32,6 @@ OptimizationTracker, Result, filter_json_field, - get_kernel_name, ) diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index a4949fad..19aa0ae9 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -35,7 +35,6 @@ OptimizationTracker, Result, filter_json_field, - get_kernel_name, ) from intelliperf.utils.process import capture_subprocess_output from intelliperf.utils.regex import generate_ecma_regex_from_list diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index c35e6f39..dc656c26 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -24,7 +24,6 @@ import json import logging -import os import sys import dspy @@ -34,7 +33,6 @@ OptimizationTracker, Result, filter_json_field, - get_kernel_name, ) From 522b6540d8fdf69ddce972250a157a99796a089d Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 04:04:42 -0600 Subject: [PATCH 14/15] Add SPDX license headers to all new Accordo files --- src/accordo/_internal/codegen.py | 3 +++ src/accordo/_internal/hip_interop.py | 3 +++ src/accordo/_internal/ipc/communication.py | 3 +++ src/accordo/config.py | 3 +++ src/accordo/exceptions.py | 3 +++ src/accordo/result.py | 3 +++ src/accordo/snapshot.py | 3 +++ src/accordo/validator.py | 3 +++ 8 files changed, 24 insertions(+) diff --git a/src/accordo/_internal/codegen.py b/src/accordo/_internal/codegen.py index fed8524a..39f95f10 100644 --- a/src/accordo/_internal/codegen.py +++ b/src/accordo/_internal/codegen.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """Code generation for Accordo C++ header files.""" import logging diff --git a/src/accordo/_internal/hip_interop.py b/src/accordo/_internal/hip_interop.py index d30c822f..0dff7cc6 100644 --- a/src/accordo/_internal/hip_interop.py +++ b/src/accordo/_internal/hip_interop.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """HIP interop functions for Accordo.""" import ctypes diff --git a/src/accordo/_internal/ipc/communication.py b/src/accordo/_internal/ipc/communication.py index 251f53f8..543758cd 100644 --- a/src/accordo/_internal/ipc/communication.py +++ b/src/accordo/_internal/ipc/communication.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """IPC communication for Accordo.""" import ctypes diff --git a/src/accordo/config.py b/src/accordo/config.py index 43efc69f..b7c6d295 100644 --- a/src/accordo/config.py +++ b/src/accordo/config.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """Configuration classes for Accordo validation.""" from dataclasses import dataclass, field diff --git a/src/accordo/exceptions.py b/src/accordo/exceptions.py index b9476fd4..5ead2700 100644 --- a/src/accordo/exceptions.py +++ b/src/accordo/exceptions.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """Custom exceptions for Accordo validation.""" diff --git a/src/accordo/result.py b/src/accordo/result.py index 9a824b66..69bd649a 100644 --- a/src/accordo/result.py +++ b/src/accordo/result.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """Result classes for Accordo validation.""" from dataclasses import dataclass diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py index 6b08a506..025f5020 100644 --- a/src/accordo/snapshot.py +++ b/src/accordo/snapshot.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """Snapshot: Represents captured kernel argument data from a binary execution.""" from dataclasses import dataclass diff --git a/src/accordo/validator.py b/src/accordo/validator.py index 13f67723..f38b29d8 100644 --- a/src/accordo/validator.py +++ b/src/accordo/validator.py @@ -1,3 +1,6 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + """AccordoValidator: Main validation class for Accordo.""" import logging From d9491681955f2deaaa6d49c2f4085af329e3bd28 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Mon, 3 Nov 2025 04:28:39 -0600 Subject: [PATCH 15/15] Re-integrate improvements from main (Unify counters report #155) - Add OptimizationTracker helper methods: is_successful(), get_best_code(), get_best_report(), get_best_metrics() - Add auto-determination of success based on improvement and speedup - Add _filter_compiler_errors() static method for cleaner compiler output - Add process crash detection in Accordo IPC with process_pid - Add enhanced timeout error messages with baseline context - Maintain all Accordo API improvements from feature branch --- src/accordo/_internal/ipc/communication.py | 20 ++- src/intelliperf/formulas/formula_base.py | 144 ++++++++++++++++++--- 2 files changed, 141 insertions(+), 23 deletions(-) diff --git a/src/accordo/_internal/ipc/communication.py b/src/accordo/_internal/ipc/communication.py index 543758cd..c5f61265 100644 --- a/src/accordo/_internal/ipc/communication.py +++ b/src/accordo/_internal/ipc/communication.py @@ -115,15 +115,31 @@ def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30, pr start_time = time.time() with open(pipe_name, "rb") as fifo: # noqa: F841 while True: + # Check if the process is still alive (crash detection, not timeout) + if process_pid is not None: + try: + os.kill(process_pid, 0) # Signal 0 just checks if process exists + except OSError: + raise RuntimeError( + f"Accordo process (PID {process_pid}) crashed or terminated during execution. " + "Check for segfaults or GPU memory access errors." + ) + if time.time() - start_time > ipc_timeout_seconds: - raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data") + timeout_msg = f"Timeout after {ipc_timeout_seconds} seconds during IPC communication" + if baseline_time_ms is not None: + timeout_msg += f" (baseline: {baseline_time_ms}ms, 2x timeout: {ipc_timeout_seconds}s). Code may be correct but too slow to be worth profiling." + raise TimeoutError(timeout_msg) try: ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) break except Exception as e: if time.time() - start_time > ipc_timeout_seconds: - raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data: {str(e)}") + timeout_msg = f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data: {str(e)}" + if baseline_time_ms is not None: + timeout_msg += f" (baseline: {baseline_time_ms}ms, 2x timeout: {ipc_timeout_seconds}s). Code may be correct but too slow to be worth profiling." + raise TimeoutError(timeout_msg) time.sleep(0.1) type_map = { diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 6c97a495..d82ca9e1 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -85,31 +85,67 @@ def __init__( # History messages for DSPy (stored as list of dicts) self.history_messages = [] + def is_successful(self) -> bool: + """Check if optimization was successful (any improvement over baseline)""" + if not self.best_step or not self.best_step.success: + return False + improvement = self.best_step.get_metric(self.primary_metric, 1.0) + speedup = self.best_step.get_metric("speedup", 1.0) + return improvement > 1.0 or speedup > 1.0 + + def get_best_code(self) -> str: + """Get the optimized code from the best step""" + if self.best_step: + return self.best_step.optimized_code + return "" + + def get_best_report(self) -> str: + """Get the optimization report from the best step""" + if self.best_step: + return self.best_step.report + return "" + + def get_best_metrics(self) -> dict: + """Get the metrics from the best step""" + if self.best_step: + return self.best_step.metrics + return {} + def add_step( self, diff: str, report: str, metrics: dict, - success: bool, + success: bool = None, request: str = "", optimized_code: str = "", ) -> OptimizationStep: - """Add step and auto-update best based on primary metric""" + """Add step and auto-calculate improvement, speedup, and success""" + # Auto-calculate speedup if time metrics are available + unopt_time = metrics.get("unoptimized_time", 0) + opt_time = metrics.get("optimized_time", 0) + if unopt_time > 0 and opt_time > 0: + metrics["speedup"] = unopt_time / opt_time + # Auto-calculate improvement if before/after metrics are configured if self.before_metric and self.after_metric: before = metrics.get(self.before_metric, 0) after = metrics.get(self.after_metric, 0) - if before != 0 and after != 0: - # Calculate improvement ratio based on whether we're maximizing or minimizing the raw metric - # For conflicts/latency (maximize=False, lower is better): improvement = before / after - # Example: 3.5 conflicts → 1.0 conflicts = 3.5 / 1.0 = 3.5x improvement - # For coalescing (maximize=True, higher is better): improvement = after / before - # Example: 50% → 75% = 75 / 50 = 1.5x improvement - improvement = before / after if not self.maximize else after / before + if before != 0: + # If maximize=False: we want to minimize raw metric (conflicts, latency), so improvement = before / after + # If maximize=True: we want to maximize raw metric (coalescing %, hit rate), so improvement = after / before + if not self.maximize: + improvement = before / after if after != 0 else 1.0 + else: + improvement = after / before if after != 0 else 1.0 metrics[self.primary_metric] = improvement - elif before != 0: - # If after is 0, set improvement to 1.0 (no change) - metrics[self.primary_metric] = 1.0 + + # Auto-determine success if not explicitly provided + if success is None: + improvement = metrics.get(self.primary_metric, 1.0) + speedup = metrics.get("speedup", 1.0) + # Success = metric improved (> 1.0) AND runtime didn't regress (>= 1.0) + success = (improvement > 1.0) and (speedup >= 1.0) step = OptimizationStep( iteration=self.current_iteration, @@ -117,6 +153,7 @@ def add_step( report=report, metrics=metrics, success=success, + optimized_code=optimized_code, ) self.steps.append(step) self.current_iteration += 1 @@ -151,15 +188,24 @@ def add_step( self.history_messages.append(history_entry) - # Auto-update best - if self.best_step is None: - self.best_step = step - else: - new_val = step.get_metric(self.primary_metric) - cur_val = self.best_step.get_metric(self.primary_metric) - - if (self.maximize and new_val > cur_val) or (not self.maximize and new_val < cur_val): + # Auto-update best based on primary metric (only for successful steps) + if success: + if self.best_step is None or not self.best_step.success: + # First successful step, or replacing a failed step self.best_step = step + else: + new_val = step.get_metric(self.primary_metric) + cur_val = self.best_step.get_metric(self.primary_metric) + + # With proper improvement calculation, we always want higher improvement values + if new_val > cur_val: + self.best_step = step + elif new_val == cur_val: + # Tie-breaker: prefer higher speedup (better runtime performance) + new_speedup = step.get_metric("speedup") + cur_speedup = self.best_step.get_metric("speedup") + if new_speedup > cur_speedup: + self.best_step = step return step @@ -395,11 +441,67 @@ def build(self, validate_build_result=True): if success: return Result(success=success, asset={"log": result}) else: + # Filter compiler log to remove noise and keep only errors + filtered_log = self._filter_compiler_errors(result) return Result( success=success, - error_report="The application contains compiler errors. Here is the compiler log: " + result, + error_report="The application contains compiler errors. Here is the compiler log:\n" + filtered_log, ) + @staticmethod + def _filter_compiler_errors(compiler_log: str, max_lines: int = 50) -> str: + """Filter compiler log to show only errors and a summary, not all warnings.""" + lines = compiler_log.split("\n") + + errors = [] + notes = [] + gmake_errors = [] + warning_count = 0 + + for line in lines: + if ": error:" in line: + errors.append(line) + elif ": note:" in line: + notes.append(line) + elif line.strip().startswith("gmake") and ("***" in line or "Error" in line): + gmake_errors.append(line) + elif ": warning:" in line: + warning_count += 1 + + # Build filtered output + filtered = [] + + if warning_count > 0: + filtered.append(f"[{warning_count} warnings omitted - only showing errors]\n") + + if errors: + filtered.append("=== COMPILATION ERRORS ===") + for error in errors[:max_lines]: # Limit errors too + filtered.append(error) + if len(errors) > max_lines: + filtered.append(f"... and {len(errors) - max_lines} more errors") + + if notes: + filtered.append("\n=== NOTES ===") + for note in notes[:max_lines]: # Limit notes too + filtered.append(note) + + if gmake_errors: + filtered.append("\n=== BUILD FAILED ===") + for gmake_error in gmake_errors[-5:]: # Last 5 gmake errors + filtered.append(gmake_error) + + if not filtered: + # No errors found, maybe it's a different kind of failure + # Return first and last few lines + filtered.append("=== BUILD OUTPUT (TRUNCATED) ===") + filtered.extend(lines[:10]) + if len(lines) > 20: + filtered.append(f"\n... {len(lines) - 20} lines omitted ...\n") + filtered.extend(lines[-10:]) + + return "\n".join(filtered) + # ---------------------------------------------------- # Required methods to be implemented by child classes # ----------------------------------------------------