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..a1613ccd 100644 --- a/src/accordo/__init__.py +++ b/src/accordo/__init__.py @@ -1,29 +1,83 @@ -################################################################################ -# 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 + - 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 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 ValidationConfig, KernelArg, AccordoValidator + >>> config = ValidationConfig( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... 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 + ... ) -# 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 AccordoValidator -"""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"] +# Public API +__all__ = [ + # Config + "ValidationConfig", + "KernelArg", + # Validator + "AccordoValidator", + # Snapshot + "Snapshot", + # Results + "ValidationResult", + "ArrayMismatch", + # Exceptions + "AccordoError", + "AccordoBuildError", + "AccordoTimeoutError", + "AccordoProcessError", + "AccordoValidationError", +] diff --git a/src/accordo/_internal/__init__.py b/src/accordo/_internal/__init__.py new file mode 100644 index 00000000..7a0b6b33 --- /dev/null +++ b/src/accordo/_internal/__init__.py @@ -0,0 +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 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..c8ef421d --- /dev/null +++ b/src/accordo/_internal/ipc/__init__.py @@ -0,0 +1 @@ +"""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..f38b29d8 --- /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 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: + # 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..9cfbb1f6 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -36,13 +36,7 @@ class AtomicContentionOptimization(dspy.Signature): - """Optimize GPU kernel code to reduce atomic contention and improve performance. - - YOU HAVE TERMINAL ACCESS: You can use the execute_terminal_command tool to explore the codebase, - read files, search for patterns, and DIRECTLY edit the kernel files. Instead of returning full code, - you can use commands like 'pwd', 'ls', 'find', 'grep', 'cat', 'sed', etc. to understand and modify files. - Your commands execute in a sandboxed project directory. Use 'pwd' to see where you are. - """ + """Optimize GPU kernel code to reduce atomic contention and improve performance.""" kernel_code = dspy.InputField(desc="The current kernel source code that needs atomic contention optimization.") @@ -259,11 +253,11 @@ def __init__( self.bottleneck_report = None self.current_summary = None self.previous_source_code = None + self.success = False # Initialize optimization tracker - # Atomic contention optimization: minimize latency (lower is better) - # Automatically calculates latency_improvement = unoptimized_lat / optimized_lat - # Higher improvement ratio is better (e.g., 100 / 50 = 2.0x improvement) + # Atomic contention optimization maximizes latency reduction (minimize atomic_lat) + # Automatically calculates latency_improvement from unoptimized_lat / optimized_lat self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="latency_improvement", @@ -276,6 +270,13 @@ def __init__( self.baseline_atomic_latency = None self.baseline_time_ms = None + # Track best optimization across iterations + self.best_speedup = 1.0 # Start at 1.0x (no speedup) + self.best_latency_improvement = 1.0 # Start at 1.0x (no improvement) + self.best_kernel_code = "" + self.best_iteration_report = "" + self.best_optimization_results = None + def profile_pass(self) -> Result: """ Profile the application using guided-tuning and collect atomic contention data @@ -357,8 +358,8 @@ def build_pass(self, validate_build_result=True) -> Result: "speedup": 0.0, "unoptimized_time": 0, "optimized_time": 0, - "unoptimized_lat": self.baseline_atomic_latency or 0, - "optimized_lat": self.baseline_atomic_latency or 0, + "unoptimized_latency": self.baseline_atomic_latency or 0, + "optimized_latency": self.baseline_atomic_latency or 0, }, success=False, request=f"Optimize atomic contention in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", @@ -394,8 +395,7 @@ def optimize_pass( system_prompt = ( "You are a skilled GPU HIP programmer specializing in optimizing kernels " - "to reduce atomic contention and improve performance. " - "You have terminal access to explore the codebase and edit files directly." + "to reduce atomic contention and improve performance." ) # Get LLM instance (initialized once per formula) @@ -525,9 +525,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={ @@ -584,18 +585,15 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) failed_code = f.read() error_report = f"Correctness Validation Failed: {result.error_report}" - # For correctness failures, we don't have valid metrics since the code didn't run correctly - # Set latency_improvement to 1.0 (no change) to indicate no progress self.optimization_tracker.add_step( diff=diff, report=error_report, metrics={ - "speedup": 1.0, # No change - "unoptimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "optimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "unoptimized_lat": self.baseline_atomic_latency or 0, - "optimized_lat": self.baseline_atomic_latency or 0, # Same as baseline (no improvement) - "latency_improvement": 1.0, # Explicitly set to 1.0 (no change) + "speedup": 0.0, + "unoptimized_time": 0, + "optimized_time": 0, + "unoptimized_latency": self.baseline_atomic_latency or 0, + "optimized_latency": self.baseline_atomic_latency or 0, }, success=False, request=f"Optimize atomic contention in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", @@ -643,51 +641,28 @@ def performance_validation_pass(self) -> Result: optimized_time = optimized_results[0]["durations"]["ns"] optimized_metric = optimized_results[0][field][subfield] - # Add step to optimization tracker (always - for learning) - # Tracker will automatically calculate: speedup, latency_improvement, and success - diff = self.compute_diff(self.current_kernel_files) - - # Read the optimized code to store in history - with open(self.current_kernel_files[0], "r") as f: - optimized_code = f.read() - - step = self.optimization_tracker.add_step( - diff=diff, - report="", # Will build report after calculations - metrics={ - "unoptimized_time": unoptimized_time, - "optimized_time": optimized_time, - "unoptimized_lat": unoptimized_metric, - "optimized_lat": optimized_metric, - }, - request=f"Optimize atomic contention in kernel {self.current_kernel_signature}", - optimized_code=optimized_code, - ) + success = optimized_metric < unoptimized_metric + speedup = unoptimized_time / optimized_time + metric_improvement = unoptimized_metric / optimized_metric if optimized_metric != 0 else 1 - # Query calculated values from tracker - speedup = step.metrics.get("speedup", 1.0) - latency_improvement = step.metrics.get("latency_improvement", 1.0) - - # Calculate percentage improvement for readability + # Calculate cycle latency improvement percentage cycle_latency_improvement = ( (unoptimized_metric - optimized_metric) / unoptimized_metric * 100 if unoptimized_metric > 0 else 0 ) - - # Build report using tracker's calculated values self.optimization_report = "" # Format the atomic contention improvement message - if latency_improvement > 1: + if metric_improvement > 1: self.optimization_report += ( f"Atomic Contention Reduction: Successfully reduced atomic contention by " - f"{latency_improvement:.2f}x. " + f"{metric_improvement:.2f}x. " f"Average atomic latency improved from {unoptimized_metric:.0f} to {optimized_metric:.0f} cycles " f"({cycle_latency_improvement:.1f}% reduction - lower latency means less contention). " ) else: self.optimization_report += ( f"Atomic Contention Increase: Atomic contention increased by " - f"{1 / latency_improvement:.2f}x. " + f"{1 / metric_improvement:.2f}x. " f"Average atomic latency worsened from {unoptimized_metric:.0f} to {optimized_metric:.0f} cycles " f"({abs(cycle_latency_improvement):.1f}% increase - higher latency means more contention). " ) @@ -706,20 +681,17 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Update the step's report with the built message - step.report = self.optimization_report - - # Log performance validation results (using tracker's calculated values) + # Log performance validation results (always, even if failed) self.get_logger().record( "performance_validation_complete", { - "success": step.success, + "success": success, "unoptimized_time_ns": unoptimized_time, "optimized_time_ns": optimized_time, "unoptimized_metric": unoptimized_metric, "optimized_metric": optimized_metric, "speedup": speedup, - "metric_improvement": latency_improvement, + "metric_improvement": metric_improvement, "cycle_latency_improvement": cycle_latency_improvement, "optimization_report": self.optimization_report, }, @@ -727,6 +699,42 @@ def performance_validation_pass(self) -> Result: logging.info(self.optimization_report) + # Add step to optimization tracker (always - for learning) + # Tracker will automatically calculate latency_improvement from before/after values + diff = self.compute_diff(self.current_kernel_files) + + # Read the optimized code to store in history + with open(self.current_kernel_files[0], "r") as f: + optimized_code = f.read() + + self.optimization_tracker.add_step( + diff=diff, + report=self.optimization_report, + metrics={ + "speedup": speedup, + "unoptimized_time": unoptimized_time, + "optimized_time": optimized_time, + "unoptimized_lat": unoptimized_metric, + "optimized_lat": optimized_metric, + }, + success=success and speedup >= 1, + request=f"Optimize atomic contention in kernel {self.current_kernel_signature}", + optimized_code=optimized_code, + ) + + # Update best if this iteration improved both speedup and latency + is_better = speedup > self.best_speedup and metric_improvement > self.best_latency_improvement + + if is_better: + self.best_speedup = speedup + self.best_latency_improvement = metric_improvement + self.best_iteration_report = self.optimization_report + self.best_optimization_results = self._optimization_results + self.best_kernel_code = optimized_code + # Mark as successful if we achieved any improvement + if metric_improvement > 1.0 or speedup > 1.0: + self.success = True + self.current_summary = self.optimization_report # Always return False to continue through all iterations @@ -736,43 +744,36 @@ def write_results(self, output_file: str = None): """ Writes the results to the output file using the best optimization attempt. """ - # Restore best code from tracker - best_code = self.optimization_tracker.get_best_code() - if best_code: - for file in self.current_kernel_files: - with open(file, "w") as f: - f.write(best_code) + # Restore best results for output + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report - # Extract metrics from best optimization step - metrics = self.optimization_tracker.get_best_metrics() + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) - # Determine metric_after: use best if available, otherwise null to indicate no improvement - metric_after = None - time_after_ms = None - if metrics and "optimized_lat" in metrics: - # Only use the measured value if it actually improved - if metrics["optimized_lat"] < self.baseline_atomic_latency: - metric_after = metrics["optimized_lat"] - time_after_ms = metrics.get("optimized_time", 0) / 1e6 + # Extract metrics from best optimization step + best_step = self.optimization_tracker.to_dict().get("best_step", {}) + metrics = best_step.get("metrics", {}) # Build structured metric fields metric_fields = { "kernel_name": self.current_kernel, "metric": "atomic_lat_cycles", # The counter we're optimizing "metric_name": "Atomic Latency", # Human-readable name - "metric_before": self.baseline_atomic_latency, - "metric_after": metric_after, # None if no optimization succeeded - "time_before_ms": self.baseline_time_ms if self.baseline_time_ms else 0.0, - "time_after_ms": time_after_ms, # None if no optimization succeeded + "metric_before": metrics.get("unoptimized_lat", self.baseline_atomic_latency), + "metric_after": metrics.get("optimized_lat", self.baseline_atomic_latency), + "time_before_ms": metrics.get("unoptimized_time", 0) / 1e6, # Convert ns to ms + "time_after_ms": metrics.get("optimized_time", 0) / 1e6, # Convert ns to ms } # Include optimization history in results super().write_results( output_file=output_file, additional_results={ - "formula": "atomicContention", - "success": self.optimization_tracker.is_successful(), "optimization_history": self.optimization_tracker.to_dict(), + "formula": "atomicContention", + "success": self.success, **metric_fields, }, ) diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index fed70ec7..19aa0ae9 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -266,11 +266,11 @@ def __init__( self.bottleneck_report = None self.current_summary = None self.previous_source_code = None + self.success = False # Initialize optimization tracker - # Bank conflict optimization: minimize conflicts (lower is better) - # Automatically calculates conflict_improvement = unoptimized_conflicts / optimized_conflicts - # Higher improvement ratio is better (e.g., 3.5 / 0.5 = 7.0x improvement) + # 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", @@ -283,6 +283,13 @@ def __init__( self.baseline_bank_conflicts = None self.baseline_time_ms = None + # Track best optimization across iterations + self.best_speedup = 1.0 # Start at 1.0x (no speedup) + self.best_conflict_improvement = 1.0 # Start at 1.0x (no improvement) + self.best_kernel_code = "" + self.best_iteration_report = "" + self.best_optimization_results = None + def build_pass(self, validate_build_result=True) -> Result: """ Build the application and store the summary. @@ -312,22 +319,21 @@ def build_pass(self, validate_build_result=True) -> Result: with open(self.current_kernel_files[0], "r") as f: failed_code = f.read() - error_report = f"Compilation Failed: {result.error_report}" - self.optimization_tracker.add_step( - diff=diff, - report=error_report, - metrics={ - "speedup": 1.0, # No change - "unoptimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "optimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "unoptimized_conflicts": self.baseline_bank_conflicts or 0, - "optimized_conflicts": self.baseline_bank_conflicts or 0, - "conflict_improvement": 1.0, # Explicitly set to 1.0 (no change) - }, - success=False, - request=f"Optimize bank conflicts in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", - optimized_code=failed_code, - ) + error_report = f"Compilation Failed: {result.error_report}" + self.optimization_tracker.add_step( + diff=diff, + report=error_report, + metrics={ + "speedup": 0.0, + "unoptimized_time": 0, + "optimized_time": 0, + "unoptimized_conflicts": self.baseline_bank_conflicts or 0, + "optimized_conflicts": self.baseline_bank_conflicts or 0, + }, + success=False, + request=f"Optimize bank conflicts in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", + optimized_code=failed_code, + ) else: self.get_logger().record( "build_pass_success", @@ -607,9 +613,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={ @@ -664,24 +670,21 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) with open(self.current_kernel_files[0], "r") as f: failed_code = f.read() - error_report = f"Correctness Validation Failed: {result.error_report}" - # For correctness failures, we don't have valid metrics since the code didn't run correctly - # Set conflict_improvement to 1.0 (no change) to indicate no progress - self.optimization_tracker.add_step( - diff=diff, - report=error_report, - metrics={ - "speedup": 1.0, # No change - "unoptimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "optimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "unoptimized_conflicts": self.baseline_bank_conflicts or 0, - "optimized_conflicts": self.baseline_bank_conflicts or 0, # Same as baseline (no improvement) - "conflict_improvement": 1.0, # Explicitly set to 1.0 (no change) - }, - success=False, - request=f"Optimize bank conflicts in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", - optimized_code=failed_code, - ) + error_report = f"Correctness Validation Failed: {result.error_report}" + self.optimization_tracker.add_step( + diff=diff, + report=error_report, + metrics={ + "speedup": 0.0, + "unoptimized_time": 0, + "optimized_time": 0, + "unoptimized_conflicts": self.baseline_bank_conflicts or 0, + "optimized_conflicts": self.baseline_bank_conflicts or 0, + }, + success=False, + request=f"Optimize bank conflicts in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", + optimized_code=failed_code, + ) else: self.get_logger().record( "correctness_validation_success", @@ -715,79 +718,50 @@ def performance_validation_pass(self) -> Result: optimized_time = optimized_results[0]["durations"]["ns"] optimized_conflicts = optimized_results[0]["lds"]["bc"] - # Add step to optimization tracker (always - for learning) - # Tracker will automatically calculate: speedup, conflict_improvement, and success - diff = self.compute_diff(self.current_kernel_files) - - # Read the optimized code to store in history - with open(self.current_kernel_files[0], "r") as f: - optimized_code = f.read() - - step = self.optimization_tracker.add_step( - diff=diff, - report="", # Will build report after calculations - metrics={ - "unoptimized_time": unoptimized_time, - "optimized_time": optimized_time, - "unoptimized_conflicts": unoptimized_conflicts, - "optimized_conflicts": optimized_conflicts, - }, - request=f"Optimize bank conflicts in kernel {self.current_kernel_signature}", - optimized_code=optimized_code, - ) - - # Query calculated values from tracker - speedup = step.metrics.get("speedup", 1.0) - conflict_improvement = step.metrics.get("conflict_improvement", 1.0) - - # Calculate percentage improvement for readability + success = optimized_conflicts < unoptimized_conflicts + speedup = unoptimized_time / optimized_time + conflict_improvement = unoptimized_conflicts / optimized_conflicts if optimized_conflicts != 0 else 1 conflict_improvement_percentage = ( - (unoptimized_conflicts - optimized_conflicts) / unoptimized_conflicts * 100 - if unoptimized_conflicts > 0 - else 0 - ) + (unoptimized_conflicts - optimized_conflicts) / unoptimized_conflicts if unoptimized_conflicts != 0 else 0 + ) * 100 - # Build report using tracker's calculated values self.optimization_report = "" # Format the conflict improvement message - if conflict_improvement > 1: + if conflict_improvement_percentage > 1: self.optimization_report += ( f"Bank Conflict Reduction: Successfully reduced shared memory bank conflicts by " - f"{conflict_improvement:.2f}x. " + f"{conflict_improvement_percentage:.1f}%. " f"Conflict ratio improved from {unoptimized_conflicts:.1f} to {optimized_conflicts:.1f} " - f"({conflict_improvement_percentage:.1f}% reduction - lower values mean fewer conflicts). " + f"(lower values indicate fewer conflicts and better performance). " ) else: self.optimization_report += ( f"Bank Conflict Increase: Bank conflicts increased by " - f"{1 / conflict_improvement:.2f}x. " + f"{abs(conflict_improvement_percentage):.1f}%. " f"Conflict ratio worsened from {unoptimized_conflicts:.1f} to {optimized_conflicts:.1f} " - f"({abs(conflict_improvement_percentage):.1f}% increase - higher values mean more conflicts). " + f"(higher values indicate more conflicts and worse performance). " ) # Format the performance improvement message if speedup > 1: self.optimization_report += ( f"Performance Gain: Achieved {speedup:.2f}x speedup with execution time " - f"reduced from {unoptimized_time / 1e6:.2f}ms to {optimized_time / 1e6:.2f}ms " + f"reduced from {unoptimized_time / 1_000_000:.2f}ms to {optimized_time / 1_000_000:.2f}ms " f"({(speedup - 1) * 100:.1f}% faster)." ) else: self.optimization_report += ( f"Performance Loss: Experienced {1 / speedup:.2f}x slowdown with execution time " - f"increased from {unoptimized_time / 1e6:.2f}ms to {optimized_time / 1e6:.2f}ms " + f"increased from {unoptimized_time / 1_000_000:.2f}ms to {optimized_time / 1_000_000:.2f}ms " f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Update the step's report with the built message - step.report = self.optimization_report - - # Log performance validation results (using tracker's calculated values) + # Log performance validation results (always, even if failed) self.get_logger().record( "performance_validation_complete", { - "success": step.success, + "success": success, "unoptimized_time_ns": unoptimized_time, "optimized_time_ns": optimized_time, "unoptimized_conflicts": unoptimized_conflicts, @@ -801,6 +775,42 @@ def performance_validation_pass(self) -> Result: logging.info(self.optimization_report) + # Add step to optimization tracker (always - for learning) + # Tracker will automatically calculate conflict_improvement from before/after values + diff = self.compute_diff(self.current_kernel_files) + + # Read the optimized code to store in history + with open(self.current_kernel_files[0], "r") as f: + optimized_code = f.read() + + self.optimization_tracker.add_step( + diff=diff, + report=self.optimization_report, + metrics={ + "speedup": speedup, + "unoptimized_time": unoptimized_time, + "optimized_time": optimized_time, + "unoptimized_conflicts": unoptimized_conflicts, + "optimized_conflicts": optimized_conflicts, + }, + success=success and speedup >= 1, + request=f"Optimize bank conflicts in kernel {self.current_kernel_signature}", + optimized_code=optimized_code, + ) + + # Update best if this iteration improved both speedup and conflict reduction + is_better = speedup > self.best_speedup and conflict_improvement > self.best_conflict_improvement + + if is_better: + self.best_speedup = speedup + self.best_conflict_improvement = conflict_improvement + self.best_iteration_report = self.optimization_report + self.best_optimization_results = self._optimization_results + self.best_kernel_code = optimized_code + # Mark as successful if we achieved any improvement + if conflict_improvement > 1.0 or speedup > 1.0: + self.success = True + self.current_summary = self.optimization_report # Always return False to continue through all iterations @@ -810,43 +820,36 @@ def write_results(self, output_file: str = None): """ Writes the results to the output file using the best optimization attempt. """ - # Restore best code from tracker - best_code = self.optimization_tracker.get_best_code() - if best_code: - for file in self.current_kernel_files: - with open(file, "w") as f: - f.write(best_code) + # Restore best results for output + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report - # Extract metrics from best optimization step - metrics = self.optimization_tracker.get_best_metrics() + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) - # Determine metric_after: use best if available, otherwise null to indicate no improvement - metric_after = None - time_after_ms = None - if metrics and "optimized_conflicts" in metrics: - # Only use the measured value if it actually improved - if metrics["optimized_conflicts"] < self.baseline_bank_conflicts: - metric_after = metrics["optimized_conflicts"] - time_after_ms = metrics.get("optimized_time", 0) / 1e6 + # Extract metrics from best optimization step + best_step = self.optimization_tracker.to_dict().get("best_step", {}) + metrics = best_step.get("metrics", {}) # Build structured metric fields metric_fields = { "kernel_name": self.current_kernel, "metric": "lds_bank_conflict", # The counter we're optimizing "metric_name": "LDS Bank Conflicts", # Human-readable name - "metric_before": self.baseline_bank_conflicts, - "metric_after": metric_after, # None if no optimization succeeded - "time_before_ms": self.baseline_time_ms if self.baseline_time_ms else 0.0, - "time_after_ms": time_after_ms, # None if no optimization succeeded + "metric_before": metrics.get("unoptimized_conflicts", self.baseline_bank_conflicts), + "metric_after": metrics.get("optimized_conflicts", self.baseline_bank_conflicts), + "time_before_ms": metrics.get("unoptimized_time", 0) / 1e6, # Convert ns to ms + "time_after_ms": metrics.get("optimized_time", 0) / 1e6, # Convert ns to ms } # Include optimization history in results super().write_results( output_file=output_file, additional_results={ - "formula": "bankConflict", - "success": self.optimization_tracker.is_successful(), "optimization_history": self.optimization_tracker.to_dict(), + "formula": "bankConflict", + "success": self.success, **metric_fields, }, ) diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 8af86207..d82ca9e1 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 AccordoValidator, KernelArg, ValidationConfig 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 @@ -58,7 +54,6 @@ class OptimizationStep: metrics: dict success: bool timestamp: float = field(default_factory=time.time) - optimized_code: str = "" # Store the optimized code for this step def get_metric(self, key: str, default=0.0): """Helper to safely get metrics""" @@ -293,6 +288,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( @@ -387,70 +386,6 @@ def get_llm(self, system_prompt: str): logging.debug(f"Initialized LLM once for formula: {self.model} via {self.provider}") return self._llm - 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 _parse_kernel_signature(self, kernel_signature: str): """ Parses a kernel signature to extract the kernel name and its arguments. @@ -596,7 +531,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 +545,68 @@ 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 = [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, ) - process_pid = process.pid - logging.debug(f"Launched {label} process with PID: {process_pid}") + 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, 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): @@ -795,6 +676,96 @@ 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 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: @@ -858,9 +829,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: diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index fc38afb7..dc656c26 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -250,15 +250,15 @@ def __init__( self.bottleneck_report = None self.current_summary = None self.previous_source_code = None + self.success = False # Initialize optimization tracker - # Memory access optimization: maximize coalescing (higher % is better) - # Automatically calculates coal_improvement = optimized_coal / unoptimized_coal - # Higher improvement ratio is better (e.g., 90% / 50% = 1.8x improvement) + # Memory access optimization maximizes coalescing improvement + # Automatically calculates coal_improvement from unoptimized_coal / optimized_coal self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="coal_improvement", - maximize=True, # True = maximize raw metric (coalescing %) + maximize=True, before_metric="unoptimized_coal", after_metric="optimized_coal", ) @@ -267,6 +267,13 @@ def __init__( self.baseline_coalesced_pct = None self.baseline_time_ms = None + # Track best optimization across iterations + self.best_speedup = 1.0 # Start at 1.0x (no speedup) + self.best_coal_improvement = 1.0 # Start at 1.0x (no improvement) + self.best_kernel_code = "" + self.best_iteration_report = "" + self.best_optimization_results = None + def build_pass(self, validate_build_result=True) -> Result: """ Build the application and store the summary. @@ -296,22 +303,21 @@ def build_pass(self, validate_build_result=True) -> Result: with open(self.current_kernel_files[0], "r") as f: failed_code = f.read() - error_report = f"Compilation Failed: {result.error_report}" - self.optimization_tracker.add_step( - diff=diff, - report=error_report, - metrics={ - "speedup": 1.0, # No change - "unoptimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "optimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "unoptimized_coal": self.baseline_coalesced_pct or 0, - "optimized_coal": self.baseline_coalesced_pct or 0, - "coal_improvement": 1.0, # Explicitly set to 1.0 (no change) - }, - success=False, - request=f"Optimize memory coalescing in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", - optimized_code=failed_code, - ) + error_report = f"Compilation Failed: {result.error_report}" + self.optimization_tracker.add_step( + diff=diff, + report=error_report, + metrics={ + "speedup": 0.0, + "unoptimized_time": 0, + "optimized_time": 0, + "unoptimized_coal": self.baseline_coalesced_pct or 0, + "optimized_coal": self.baseline_coalesced_pct or 0, + }, + success=False, + request=f"Optimize memory coalescing in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", + optimized_code=failed_code, + ) else: self.get_logger().record( "build_pass_success", @@ -514,9 +520,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={ @@ -572,24 +578,21 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) with open(self.current_kernel_files[0], "r") as f: failed_code = f.read() - error_report = f"Correctness Validation Failed: {result.error_report}" - # For correctness failures, we don't have valid metrics since the code didn't run correctly - # Set coal_improvement to 1.0 (no change) to indicate no progress - self.optimization_tracker.add_step( - diff=diff, - report=error_report, - metrics={ - "speedup": 1.0, # No change - "unoptimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "optimized_time": self.baseline_time_ms * 1e6 if self.baseline_time_ms else 0, - "unoptimized_coal": self.baseline_coalesced_pct or 0, - "optimized_coal": self.baseline_coalesced_pct or 0, # Same as baseline (no improvement) - "coal_improvement": 1.0, # Explicitly set to 1.0 (no change) - }, - success=False, - request=f"Optimize memory coalescing in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", - optimized_code=failed_code, - ) + error_report = f"Correctness Validation Failed: {result.error_report}" + self.optimization_tracker.add_step( + diff=diff, + report=error_report, + metrics={ + "speedup": 0.0, + "unoptimized_time": 0, + "optimized_time": 0, + "unoptimized_coal": self.baseline_coalesced_pct or 0, + "optimized_coal": self.baseline_coalesced_pct or 0, + }, + success=False, + request=f"Optimize memory coalescing in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", + optimized_code=failed_code, + ) else: self.get_logger().record( "correctness_validation_success", @@ -623,54 +626,26 @@ def performance_validation_pass(self) -> Result: optimized_time = optimized_results[0]["durations"]["ns"] optimized_coal = optimized_results[0]["l1"]["coal"] - # Add step to optimization tracker (always - for learning) - # Tracker will automatically calculate: speedup, coal_improvement, and success - diff = self.compute_diff(self.current_kernel_files) - - # Read the optimized code to store in history - with open(self.current_kernel_files[0], "r") as f: - optimized_code = f.read() + success = optimized_coal > unoptimized_coal + speedup = unoptimized_time / optimized_time + coal_improvement = optimized_coal / unoptimized_coal if optimized_coal != 0 else 1 - step = self.optimization_tracker.add_step( - diff=diff, - report="", # Will build report after calculations - metrics={ - "unoptimized_time": unoptimized_time, - "optimized_time": optimized_time, - "unoptimized_coal": unoptimized_coal, - "optimized_coal": optimized_coal, - }, - request=f"Optimize memory access pattern in kernel {self.current_kernel_signature}", - optimized_code=optimized_code, - ) - - # Query calculated values from tracker - speedup = step.metrics.get("speedup", 1.0) - coal_improvement = step.metrics.get("coal_improvement", 1.0) - - # Build report using tracker's calculated values self.optimization_report = "" # Format the memory coalescing improvement message if coal_improvement > 1: - coal_pct_improvement = ( - ((optimized_coal - unoptimized_coal) / unoptimized_coal * 100) if unoptimized_coal > 0 else 0 - ) self.optimization_report += ( f"Memory Coalescing Improvement: Successfully improved memory access patterns by " f"{coal_improvement:.2f}x. " f"Coalescing efficiency increased from {unoptimized_coal:.1f}% to {optimized_coal:.1f}% " - f"({coal_pct_improvement:.1f}% increase - higher percentages indicate more efficient memory access). " + f"(higher percentages indicate more efficient memory access patterns). " ) else: - coal_pct_decrease = ( - ((unoptimized_coal - optimized_coal) / unoptimized_coal * 100) if unoptimized_coal > 0 else 0 - ) self.optimization_report += ( f"Memory Coalescing Degradation: Memory access patterns worsened by " f"{1 / coal_improvement:.2f}x. " f"Coalescing efficiency decreased from {unoptimized_coal:.1f}% to {optimized_coal:.1f}% " - f"({coal_pct_decrease:.1f}% decrease - lower percentages indicate less efficient memory access). " + f"(lower percentages indicate less efficient memory access patterns). " ) # Format the performance improvement message @@ -687,14 +662,11 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Update the step's report with the built message - step.report = self.optimization_report - - # Log performance validation results (using tracker's calculated values) + # Log performance validation results (always, even if failed) self.get_logger().record( "performance_validation_complete", { - "success": step.success, + "success": success, "unoptimized_time_ns": unoptimized_time, "optimized_time_ns": optimized_time, "unoptimized_coal": unoptimized_coal, @@ -707,6 +679,42 @@ def performance_validation_pass(self) -> Result: logging.info(self.optimization_report) + # Add step to optimization tracker (always - for learning) + # Tracker will automatically calculate coal_improvement from before/after values + diff = self.compute_diff(self.current_kernel_files) + + # Read the optimized code to store in history + with open(self.current_kernel_files[0], "r") as f: + optimized_code = f.read() + + self.optimization_tracker.add_step( + diff=diff, + report=self.optimization_report, + metrics={ + "speedup": speedup, + "unoptimized_time": unoptimized_time, + "optimized_time": optimized_time, + "unoptimized_coal": unoptimized_coal, + "optimized_coal": optimized_coal, + }, + success=success and speedup >= 1, + request=f"Optimize memory access pattern in kernel {self.current_kernel_signature}", + optimized_code=optimized_code, + ) + + # Update best if this iteration improved both speedup and coalescing + is_better = speedup > self.best_speedup and coal_improvement > self.best_coal_improvement + + if is_better: + self.best_speedup = speedup + self.best_coal_improvement = coal_improvement + self.best_iteration_report = self.optimization_report + self.best_optimization_results = self._optimization_results + self.best_kernel_code = optimized_code + # Mark as successful if we achieved any improvement + if coal_improvement > 1.0 or speedup > 1.0: + self.success = True + self.current_summary = self.optimization_report # Always return False to continue through all iterations @@ -716,43 +724,36 @@ def write_results(self, output_file: str = None): """ Writes the results to the output file using the best optimization attempt. """ - # Restore best code from tracker - best_code = self.optimization_tracker.get_best_code() - if best_code: - for file in self.current_kernel_files: - with open(file, "w") as f: - f.write(best_code) + # Restore best results for output + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report - # Extract metrics from best optimization step - metrics = self.optimization_tracker.get_best_metrics() + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) - # Determine metric_after: use best if available, otherwise null to indicate no improvement - metric_after = None - time_after_ms = None - if metrics and "optimized_coal" in metrics: - # Only use the measured value if it actually improved (maximize=True, so higher is better) - if metrics["optimized_coal"] > self.baseline_coalesced_pct: - metric_after = metrics["optimized_coal"] - time_after_ms = metrics.get("optimized_time", 0) / 1e6 + # Extract metrics from best optimization step + best_step = self.optimization_tracker.to_dict().get("best_step", {}) + metrics = best_step.get("metrics", {}) # Build structured metric fields metric_fields = { "kernel_name": self.current_kernel, "metric": "coal_pct", # The counter we're optimizing (memory coalescing) "metric_name": "Memory Coalescing", # Human-readable name - "metric_before": self.baseline_coalesced_pct, - "metric_after": metric_after, # None if no optimization succeeded - "time_before_ms": self.baseline_time_ms if self.baseline_time_ms else 0.0, - "time_after_ms": time_after_ms, # None if no optimization succeeded + "metric_before": metrics.get("unoptimized_coal", self.baseline_coalesced_pct), + "metric_after": metrics.get("optimized_coal", self.baseline_coalesced_pct), + "time_before_ms": metrics.get("unoptimized_time", 0) / 1e6, # Convert ns to ms + "time_after_ms": metrics.get("optimized_time", 0) / 1e6, # Convert ns to ms } # Include optimization history in results super().write_results( output_file=output_file, additional_results={ - "formula": "memoryAccess", - "success": self.optimization_tracker.is_successful(), "optimization_history": self.optimization_tracker.to_dict(), + "formula": "memoryAccess", + "success": self.success, **metric_fields, }, )