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/__init__.py b/src/accordo/__init__.py index 946f447d..419a21a5 100644 --- a/src/accordo/__init__.py +++ b/src/accordo/__init__.py @@ -1,29 +1,105 @@ -################################################################################ -# MIT License +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +"""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 + - Snapshot: Captured kernel argument data from binary execution + - ValidationResult: Result of validation with detailed metrics + - ArrayMismatch: Information about array validation failures + - Accordo: Main validator class for kernel validation + - 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: +Quick Example (one-off validation): + >>> from accordo import Accordo + >>> config = Accordo.Config( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... Accordo.KernelArg(name="result", type="double*"), + ... Accordo.KernelArg(name="input", type="const double*"), + ... ], + ... tolerance=1e-6 + ... ) + >>> validator = Accordo(config) + >>> result = validator.validate( + ... reference_binary=["./app_ref"], + ... optimized_binary=["./app_opt"], + ... working_directory=".", + ... baseline_time_ms=10.0 + ... ) -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. +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) +""" -# 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. -################################################################################ +# Public API exports +from .config import KernelArg, ValidationConfig +from .exceptions import ( + AccordoBuildError, + AccordoError, + AccordoProcessError, + AccordoTimeoutError, + AccordoValidationError, +) +from .result import ArrayMismatch, ValidationResult +from .snapshot import Snapshot +from .validator import Accordo as _Accordo -"""Accordo package for validation and verification.""" +# Version +__version__ = "0.2.0" -from .python import code_gen, communicate, hip, utils -__all__ = ["communicate", "code_gen", "utils", "hip"] +# Nest all classes under Accordo namespace +class Accordo(_Accordo): + """Main Accordo validator with nested classes for clean API. + + All Accordo components are accessible as Accordo.ClassName: + - Accordo.Config (ValidationConfig) + - Accordo.KernelArg + - Accordo.Snapshot + - Accordo.Result (ValidationResult) + - Accordo.ArrayMismatch + - Accordo.Error (AccordoError) + - Accordo.BuildError (AccordoBuildError) + - Accordo.TimeoutError (AccordoTimeoutError) + - Accordo.ProcessError (AccordoProcessError) + - Accordo.ValidationError (AccordoValidationError) + """ + + # Configuration + Config = ValidationConfig + KernelArg = KernelArg + + # Data structures + Snapshot = Snapshot + Result = ValidationResult + ArrayMismatch = ArrayMismatch + + # Exceptions + Error = AccordoError + BuildError = AccordoBuildError + TimeoutError = AccordoTimeoutError + ProcessError = AccordoProcessError + ValidationError = AccordoValidationError + + +# Public API +__all__ = [ + "Accordo", +] diff --git a/src/accordo/_internal/__init__.py b/src/accordo/_internal/__init__.py new file mode 100644 index 00000000..cfc5bea3 --- /dev/null +++ b/src/accordo/_internal/__init__.py @@ -0,0 +1,3 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +"""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..39f95f10 --- /dev/null +++ b/src/accordo/_internal/codegen.py @@ -0,0 +1,53 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""Code generation for Accordo C++ header files.""" + +import logging + + +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_interop.py b/src/accordo/_internal/hip_interop.py new file mode 100644 index 00000000..0dff7cc6 --- /dev/null +++ b/src/accordo/_internal/hip_interop.py @@ -0,0 +1,109 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""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..c0c05d33 --- /dev/null +++ b/src/accordo/_internal/ipc/__init__.py @@ -0,0 +1,3 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +"""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..c5f61265 --- /dev/null +++ b/src/accordo/_internal/ipc/communication.py @@ -0,0 +1,184 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""IPC communication for Accordo.""" + +import ctypes +import logging +import os +import stat +import time + +import ml_dtypes +import numpy as np + +from ..hip_interop 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: + # 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: + 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: + 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 = { + "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..b7c6d295 --- /dev/null +++ b/src/accordo/config.py @@ -0,0 +1,109 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""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") + + Examples: + >>> 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 + + @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..5ead2700 --- /dev/null +++ b/src/accordo/exceptions.py @@ -0,0 +1,38 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""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..69bd649a --- /dev/null +++ b/src/accordo/result.py @@ -0,0 +1,103 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""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() diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py new file mode 100644 index 00000000..025f5020 --- /dev/null +++ b/src/accordo/snapshot.py @@ -0,0 +1,60 @@ +# 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 +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 = [ + "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..3d68b00d --- /dev/null +++ b/src/accordo/validator.py @@ -0,0 +1,453 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""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 ._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 +from .result import ArrayMismatch, ValidationResult +from .snapshot import Snapshot + + +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 Accordo: + """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: + # Kill the process if it timed out + try: + os.kill(process_pid, 9) + except (OSError, ProcessLookupError): + pass # Process already dead + 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/__main__.py b/src/intelliperf/__main__.py index bf2d0926..d9e5e9f0 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..65da6ed8 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 18f690a0..d188ffc4 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -525,8 +525,8 @@ 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, @@ -770,9 +770,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.optimization_tracker.is_successful(), - "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 fed70ec7..4c6eb9c8 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -607,8 +607,8 @@ 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, @@ -844,9 +844,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.optimization_tracker.is_successful(), - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) diff --git a/src/intelliperf/formulas/diagnose_only.py b/src/intelliperf/formulas/diagnose_only.py index 43da42ce..0dc2db34 100644 --- a/src/intelliperf/formulas/diagnose_only.py +++ b/src/intelliperf/formulas/diagnose_only.py @@ -70,14 +70,25 @@ def optimize_pass(self, target_kernel: str = None): def compile_pass(self): return super().compile_pass() - def correctness_validation_pass(self): + def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tolerance: float = 1e-6): """ - Validate the optimized kernel by comparing the output with the reference kernel + Validate the optimized kernel by comparing the output with the reference kernel. + + Note: diagnose_only doesn't actually optimize, so this always returns success. + The method signature matches other formulas for API consistency. + + Args: + kernel: Kernel name (unused for diagnose_only) + kernel_args: Kernel arguments (unused for diagnose_only) + accordo_absolute_tolerance: Tolerance parameter (unused for diagnose_only) Returns: - Result: Validation status + Result: Validation status (always success for diagnose_only) """ - return super().correctness_validation_pass() + # diagnose_only doesn't optimize, so validation always succeeds + from intelliperf.formulas.formula_base import Result + + return Result(success=True, asset={"log": "diagnose_only: No optimization performed, validation skipped."}) def performance_validation_pass(self): return super().performance_validation_pass() diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 8af86207..b592c9d1 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -26,7 +26,6 @@ import json import logging import os -import subprocess import sys import time from abc import abstractmethod @@ -38,13 +37,10 @@ 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 Accordo 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 @@ -349,6 +345,10 @@ def __init__( self._llm = None self._dspy_configured = False + # Accordo caching: validator and reference snapshot are created once and reused + self._accordo_validator = None + self._reference_snapshot = None + self.build() def get_logger(self) -> Logger: @@ -451,6 +451,32 @@ def find_kernel_file(self, files: list, kernel: str) -> tuple: return kernel_file, unoptimized_file_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 _parse_kernel_signature(self, kernel_signature: str): """ Parses a kernel signature to extract the kernel name and its arguments. @@ -596,7 +622,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() @@ -607,125 +636,70 @@ 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() - - # Get baseline time if available (for dynamic timeout calculation) - baseline_time_ms = getattr(self, "baseline_time_ms", None) - if baseline_time_ms is not None: - logging.debug(f"Using baseline time for dynamic timeout: {baseline_time_ms}ms") - - 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}") - - # Launch the process with Accordo and track its PID - process = subprocess.Popen( - binary_with_args, env=env, cwd=project_directory, stdout=subprocess.DEVNULL, stderr=subprocess.DEVNULL + # Create validator if not already created (and cache it) + if self._accordo_validator is None: + kernel_arg_objects = [ + Accordo.KernelArg(name=f"arg{i}", type=arg_type) for i, arg_type in enumerate(kernel_args) + ] + + config = Accordo.Config( + kernel_name=kernel, + kernel_args=kernel_arg_objects, + tolerance=accordo_absolute_tolerance, + timeout_multiplier=2.0, ) - process_pid = process.pid - logging.debug(f"Launched {label} process with PID: {process_pid}") + self._accordo_validator = Accordo(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, process_pid=process_pid, baseline_time_ms=baseline_time_ms + 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 ) - except TimeoutError as e: - logging.error(f"Timeout while getting kernel argument data for {label}: {str(e)}") - process.kill() # Kill the hung process + 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 - # Provide context-specific error message - if baseline_time_ms is not None: - error_msg = f"Optimization exceeded 2x baseline execution time for {label}: {str(e)}. Code may be correct but too slow to be worth profiling." - else: - error_msg = ( - f"Timeout while getting kernel argument data for {label}: {str(e)}. The code may have crashed." - ) + 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") - return Result( - success=False, - error_report=error_msg, - ) - except RuntimeError as e: - logging.error(f"Accordo process crashed for {label}: {str(e)}") - return Result( - success=False, - error_report=f"Accordo process crashed for {label}: {str(e)}. This usually indicates a segfault or GPU memory access error in the kernel code.", - ) - send_response(pipe_name) - - # Wait for the process to finish - process.wait(timeout=5) - 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)}") + # 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): diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index fc38afb7..3822c5c2 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -514,8 +514,8 @@ 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, @@ -750,9 +750,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.optimization_tracker.is_successful(), - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, )