From d46411ca35563fd61d882891563cdac1dcdf0b02 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Mon, 8 Sep 2025 11:56:36 -0700 Subject: [PATCH 01/14] Change Nexus repository URL to HTTPS --- pyproject.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pyproject.toml b/pyproject.toml index f066944d..6a2bdc8c 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -39,7 +39,7 @@ python3 -m pip install -r requirements.txt [tool.nexus] -git = "git@github.com:AMDResearch/nexus.git" +git = "https://github.com/AMDResearch/nexus.git" branch = "main" build_command = """ export CC=${ROCM_PATH}/bin/hipcc From 27c2c394f660153fb41c31f569291923984b1cdb Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 17 Sep 2025 02:57:21 -0700 Subject: [PATCH 02/14] Update containers and CI (#145) --- .github/workflows/ci.yml | 292 ++++++------------ .../workflows/scripts/check_test_results.sh | 41 +++ .github/workflows/scripts/ci_tests.sh | 32 +- README.md | 2 +- apptainer/build.sh | 65 +--- apptainer/intelliperf.def | 61 +--- apptainer/run.sh | 39 +-- apptainer/run_cmd.sh | 25 +- docker/build.sh | 67 +--- docker/intelliperf.Dockerfile | 19 +- docker/run.sh | 31 +- pyproject.toml | 4 +- 12 files changed, 180 insertions(+), 498 deletions(-) create mode 100755 .github/workflows/scripts/check_test_results.sh diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index b1debb00..b6a24fbe 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -1,216 +1,102 @@ -name: IntelliPerf CI +name: IntelliPerf Tests with Apptainer on: - pull_request: - branches: [ main ] push: branches: [ main ] + pull_request: + branches: [ main ] workflow_dispatch: concurrency: - group: ${{ github.workflow }}-${{ github.ref }} - cancel-in-progress: true + group: ${{ github.workflow }}-${{ github.head_ref || github.ref }} + cancel-in-progress: ${{ github.ref != 'refs/heads/main' }} jobs: - test-intelliperf: - runs-on: ubuntu-latest - env: - DIGITALOCEAN_API_URL: ${{ secrets.DIGITALOCEAN_API_URL }} - + build-apptainer-image: + runs-on: [self-hosted] + timeout-minutes: 90 + steps: - - name: Checkout code - uses: actions/checkout@v4 - - - name: Install doctl - uses: digitalocean/action-doctl@v2 - with: - token: ${{ secrets.DEV_CLOUD_KEY }} - - - name: Install jq - run: | - sudo apt-get update - sudo apt-get install -y jq - - - name: Create Droplet - id: create - env: - DIGITALOCEAN_API_URL: ${{ secrets.DIGITALOCEAN_API_URL }} - run: | - DROPLET_NAME="intelliperf-$(date +%s)" - - # Create droplet and capture JSON output - DROPLET_JSON=$(doctl compute droplet create \ - --image 188571990 \ - --size ${{ secrets.DIGITALOCEAN_SIZE }} \ - --region atl1 \ - --ssh-keys ${{ secrets.SSH_KEY_ID }} \ - "$DROPLET_NAME" \ - -o json \ - --wait 2>&1) - - # Check if droplet creation was successful - if [ $? -ne 0 ]; then - echo "❌ Failed to create droplet" - echo "Error details:" - echo "$DROPLET_JSON" - exit 1 - fi - - # Extract droplet ID and IP - DROPLET_ID=$(echo "$DROPLET_JSON" | jq -r '.[0].id') - PUBLIC_IP=$(echo "$DROPLET_JSON" | jq -r '.[0].networks.v4[] | select(.type=="public") | .ip_address') - - # Set outputs for other steps - echo "droplet_id=$DROPLET_ID" >> $GITHUB_OUTPUT - echo "public_ip=$PUBLIC_IP" >> $GITHUB_OUTPUT - - echo "βœ… Droplet created successfully!" - - - name: Setup SSH key - run: | - mkdir -p ~/.ssh - echo "${{ secrets.SSH_PRIVATE_KEY }}" > ~/.ssh/id_rsa - chmod 600 ~/.ssh/id_rsa - ssh-keyscan -H ${{ steps.create.outputs.public_ip }} >> ~/.ssh/known_hosts 2>/dev/null || true - - - name: Wait for SSH to be ready - run: | - echo "⏳ Waiting for SSH to be ready..." - for i in {1..30}; do - if ssh -o ConnectTimeout=10 -o StrictHostKeyChecking=no root@${{ steps.create.outputs.public_ip }} "echo 'SSH ready'" 2>/dev/null; then - echo "βœ… SSH is ready!" - break + - name: Checkout repository + uses: actions/checkout@v4 + + - name: Setup Apptainer + run: | + apt-get update && apt-get install -y software-properties-common + echo "deb [trusted=yes] https://ppa.launchpadcontent.net/apptainer/ppa/ubuntu $(lsb_release -cs) main" \ + | tee /etc/apt/sources.list.d/apptainer.list + apt-get update && apt-get install -y apptainer + + - name: Build IntelliPerf Apptainer container + run: | + # Create persistent Apptainer directory + mkdir -p ~/apptainer + + # Build Apptainer image from definition file (only if it doesn't exist) + if [ ! -f ~/apptainer/intelliperf-dev.sif ]; then + echo "Building new Apptainer image..." + ./apptainer/build.sh --output ~/apptainer/intelliperf-dev.sif + else + echo "Using existing Apptainer image" fi - echo "Attempt $i/30: SSH not ready yet, waiting 10 seconds..." - sleep 10 - done - - - name: Determine commit hash - id: commit_hash - run: | - if [ "${{ github.event_name }}" == "pull_request" ]; then - echo "commit_hash=${{ github.event.pull_request.head.sha }}" >> $GITHUB_OUTPUT - else - echo "commit_hash=${{ github.sha }}" >> $GITHUB_OUTPUT - fi - - - name: Install IntelliPerf and run tests - run: | - echo "πŸš€IntelliPerf installation..." - - # Setup SSH, clone repo, and install dependencies - ssh -o StrictHostKeyChecking=no root@${{ steps.create.outputs.public_ip }} " - set -e - - # Remove any stale dpkg locks - sudo rm -f /var/lib/apt/lists/lock - sudo rm -f /var/cache/apt/archives/lock - sudo rm -f /var/lib/dpkg/lock* - - # Setup SSH key for git access - mkdir -p ~/.ssh - echo '${{ secrets.SSH_PRIVATE_KEY }}' > ~/.ssh/id_rsa - chmod 600 ~/.ssh/id_rsa - ssh-keyscan -H github.com >> ~/.ssh/known_hosts - - # Set environment variables - export LLM_GATEWAY_KEY='${{ secrets.LLM_GATEWAY_KEY }}' - export ROCM_PATH=/opt/rocm - export PATH=\$ROCM_PATH/bin:\$PATH - export LD_LIBRARY_PATH=\$ROCM_PATH/lib:\$LD_LIBRARY_PATH - - # Install system dependencies - sudo apt-get update - sudo apt-get install -y python3-venv rocm-llvm-dev libzstd-dev libdwarf-dev locales git cmake - sudo locale-gen en_US.UTF-8 - - # Clone the repository - git clone git@github.com:AMDResearch/intelliperf.git - cd intelliperf - echo 'Checking out commit ${{ steps.commit_hash.outputs.commit_hash }}' - git checkout ${{ steps.commit_hash.outputs.commit_hash }} - - # Setup Python environment - python3 -m venv intelliperf_env - source intelliperf_env/bin/activate - pip install --upgrade pip - pip install -e . - pip3 install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.4 - python3 scripts/install_tool.py --all - - # Verify installation - rocminfo - pip list | grep intelliperf || echo 'IntelliPerf not found in pip list' + run-tests: + name: IntelliPerf Test + needs: build-apptainer-image + runs-on: [self-hosted] + timeout-minutes: 20 + + steps: + - name: Checkout repository + uses: actions/checkout@v4 + + - name: Run IntelliPerf Tests + env: + LLM_GATEWAY_KEY: ${{ secrets.LLM_GATEWAY_KEY }} + run: | + apptainer exec ~/apptainer/intelliperf-dev.sif bash -c " + set -e # Exit on any error + + # Set environment variables + export LLM_GATEWAY_KEY='${{ secrets.LLM_GATEWAY_KEY }}' + export ROCM_PATH=/opt/rocm + export PATH=\$ROCM_PATH/bin:\$PATH + export LD_LIBRARY_PATH=\$ROCM_PATH/lib:\$LD_LIBRARY_PATH + + # Install IntelliPerf (PyTorch ROCm and tools are already installed in container) + pip install -e . + python3 scripts/install_tool.py --all + + # IntelliPerf is installed at /root/.local/bin + export PATH=\$PATH:/root/.local/bin + + # Create results directory + mkdir -p ~/intelliperf_results + + # Run the IntelliPerf examples + ./.github/workflows/scripts/ci_tests.sh + " + - name: Check test results + run: | + ./.github/workflows/scripts/check_test_results.sh + + - name: Upload test outputs as artifact + if: always() + run: | + echo "πŸ“₯ Collecting test outputs..." + mkdir -p test_outputs - # Create results directory - mkdir -p /intelliperf_results + # Copy test results from container to local directory + cp -r ~/intelliperf_results/ ./test_outputs/ || echo "No results directory found" - # Run the IntelliPerf examples - ./.github/workflows/scripts/ci_tests.sh - " - - - name: Download test outputs - if: always() - run: | - echo "πŸ“₯ Downloading test outputs..." - mkdir -p test_outputs - scp -r -o StrictHostKeyChecking=no root@${{ steps.create.outputs.public_ip }}:/intelliperf_results/ ./test_outputs/ || echo "No results directory found" - - # Create tar artifact - tar -czf intelliperf_test_outputs.tar.gz -C test_outputs . - echo "βœ… Test outputs archived as intelliperf_test_outputs.tar.gz" - - # Print test results summary with GitHub Actions annotations - echo "πŸ“Š IntelliPerf Test Results Summary:" - - # Check each test result by parsing the success field - check_test_result() { - local file="$1" - local test_name="$2" - if [ -f "$file" ]; then - if jq -e '.success == true' "$file" >/dev/null 2>&1; then - echo "::notice::βœ… $test_name: PASSED" - return 0 - else - echo "::warning::❌ $test_name: FAILED" - return 1 - fi - else - echo "::warning::❌ $test_name: FAILED (file not found)" - return 1 - fi - } - - # Track overall success - overall_success=true - - check_test_result "./test_outputs/intelliperf_results/memory_access_output.json" "Memory Access" || overall_success=false - check_test_result "./test_outputs/intelliperf_results/bank_conflict_output.json" "Bank Conflict" || overall_success=false - check_test_result "./test_outputs/intelliperf_results/atomic_contention_output.json" "Atomic Contention" || overall_success=false - check_test_result "./test_outputs/intelliperf_results/diagnose_only_hip_uncoalesced.json" "Diagnose Only (HIP)" || overall_success=false - check_test_result "./test_outputs/intelliperf_results/diagnose_only_torch_add.json" "Diagnose Only (Torch)" || overall_success=false - check_test_result "./test_outputs/intelliperf_results/diagnose_only_triton_reduce.json" "Diagnose Only (Triton)" || overall_success=false - - echo "" - if [ "$overall_success" = true ]; then - echo "::notice::🎯 All IntelliPerf tests PASSED! βœ…" - else - echo "::warning::⚠️ Some IntelliPerf tests FAILED! ❌" - fi - - - name: Upload test outputs as artifact - if: always() - uses: actions/upload-artifact@v4 - with: - name: intelliperf-test-outputs - path: intelliperf_test_outputs.tar.gz - retention-days: 15 - - - name: Auto-destroy droplet after use - if: always() - env: - DIGITALOCEAN_API_URL: ${{ secrets.DIGITALOCEAN_API_URL }} - run: | - echo "πŸ—‘οΈ Auto-destroying droplet ${{ steps.create.outputs.droplet_id }}..." - doctl compute droplet delete ${{ steps.create.outputs.droplet_id }} --force - echo "βœ… Droplet auto-destroyed successfully!" + # Create tar artifact + tar -czf intelliperf_test_outputs.tar.gz -C test_outputs . + echo "βœ… Test outputs archived as intelliperf_test_outputs.tar.gz" + + - name: Upload artifact + if: always() + uses: actions/upload-artifact@v4 + with: + name: intelliperf-test-outputs + path: intelliperf_test_outputs.tar.gz + retention-days: 15 + diff --git a/.github/workflows/scripts/check_test_results.sh b/.github/workflows/scripts/check_test_results.sh new file mode 100755 index 00000000..cb91e0e5 --- /dev/null +++ b/.github/workflows/scripts/check_test_results.sh @@ -0,0 +1,41 @@ +#!/bin/bash +set -e + +# Check test results for success +echo 'Checking test results...' +results_dir=~/intelliperf_results + +# Function to check test result +check_test_result() { + local file="$1" + local test_name="$2" + if [ -f "$file" ]; then + if jq -e '.success == true' "$file" >/dev/null 2>&1; then + echo "βœ… $test_name: PASSED" + return 0 + else + echo "❌ $test_name: FAILED" + return 1 + fi + else + echo "❌ $test_name: FAILED (file not found)" + return 1 + fi +} + +# Track overall success +overall_success=true + +check_test_result "$results_dir/memory_access_output.json" "Memory Access" || overall_success=false +check_test_result "$results_dir/bank_conflict_output.json" "Bank Conflict" || overall_success=false +check_test_result "$results_dir/atomic_contention_output.json" "Atomic Contention" || overall_success=false +check_test_result "$results_dir/diagnose_only_hip_uncoalesced.json" "Diagnose Only (HIP)" || overall_success=false +check_test_result "$results_dir/diagnose_only_torch_add.json" "Diagnose Only (Torch)" || overall_success=false +check_test_result "$results_dir/diagnose_only_triton_reduce.json" "Diagnose Only (Triton)" || overall_success=false + +echo "" +if [ "$overall_success" = true ]; then + echo "🎯 All IntelliPerf tests PASSED! βœ…" +else + echo "⚠️ Some IntelliPerf tests FAILED! ❌" +fi diff --git a/.github/workflows/scripts/ci_tests.sh b/.github/workflows/scripts/ci_tests.sh index 02562eeb..4916838b 100755 --- a/.github/workflows/scripts/ci_tests.sh +++ b/.github/workflows/scripts/ci_tests.sh @@ -4,26 +4,32 @@ set -e # Run examples and store outputs echo 'Running IntelliPerf examples...' -mkdir -p /intelliperf_results +results_dir=~/intelliperf_results +rm -rf $results_dir +mkdir -p $results_dir + + +provider="openrouter" +model="openai/gpt-4o" # Formulas -intelliperf -vvv --project_directory=./examples --build_command="./scripts/build_examples.sh -c" --formula=memoryAccess -o /intelliperf_results/memory_access_output.json -- ./build/access_pattern/uncoalesced || true -intelliperf -vvv --project_directory=./examples --build_command="./scripts/build_examples.sh -c" --formula=bankConflict -o /intelliperf_results/bank_conflict_output.json -- ./build/bank_conflict/matrix_transpose 1024 1024 || true -intelliperf -vvv --project_directory=./examples --build_command="./scripts/build_examples.sh -c" --instrument_command="./scripts/build_examples.sh -i -c" --formula=atomicContention -o /intelliperf_results/atomic_contention_output.json -- ./build/contention/reduction || true +intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --formula=memoryAccess -o $results_dir/memory_access_output.json -- ./build/access_pattern/uncoalesced || true +intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --formula=bankConflict -o $results_dir/bank_conflict_output.json -- ./build/bank_conflict/matrix_transpose 1024 1024 || true +intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --instrument_command="./scripts/build_examples.sh -i -c" --formula=atomicContention -o $results_dir/atomic_contention_output.json -- ./build/contention/reduction || true # Diagnose Only -intelliperf -vvv --formula=diagnoseOnly -o /intelliperf_results/diagnose_only_hip_uncoalesced.json -- ./examples/build/access_pattern/uncoalesced -intelliperf -vvv --formula=diagnoseOnly -o /intelliperf_results/diagnose_only_torch_add.json -- python ./examples/torch/add.py -TRITON_DISABLE_LINE_INFO=0 intelliperf -vvv --formula=diagnoseOnly -o /intelliperf_results/diagnose_only_triton_reduce.json -- python ./examples/triton/reduce.py +intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_hip_uncoalesced.json -- ./examples/build/access_pattern/uncoalesced +intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_torch_add.json -- ./examples/torch/add.py +TRITON_DISABLE_LINE_INFO=0 intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_triton_reduce.json -- ./examples/triton/reduce.py # Display output files echo 'Memory Access Output:' -cat /intelliperf_results/memory_access_output.json || echo "File not found" +cat $results_dir/memory_access_output.json || echo "File not found" echo 'Bank Conflict Output:' -cat /intelliperf_results/bank_conflict_output.json || echo "File not found" +cat $results_dir/bank_conflict_output.json || echo "File not found" echo 'Atomic Contention Output:' -cat /intelliperf_results/atomic_contention_output.json || echo "File not found" +cat $results_dir/atomic_contention_output.json || echo "File not found" echo 'Diagnose Only Output:' -cat /intelliperf_results/diagnose_only_hip_uncoalesced.json || echo "File not found" -cat /intelliperf_results/diagnose_only_torch_add.json || echo "File not found" -cat /intelliperf_results/diagnose_only_triton_reduce.json || echo "File not found" +cat $results_dir/diagnose_only_hip_uncoalesced.json || echo "File not found" +cat $results_dir/diagnose_only_torch_add.json || echo "File not found" +cat $results_dir/diagnose_only_triton_reduce.json || echo "File not found" diff --git a/README.md b/README.md index ba8bf6c8..3095037e 100644 --- a/README.md +++ b/README.md @@ -93,7 +93,7 @@ docker run -it --rm --device=/dev/kfd --device=/dev/dri --group-add video -e LLM 1. **Clone the Repository**: ```bash - git clone git@github.com:AMDResearch/intelliperf.git + git clone https://github.com/AMDResearch/intelliperf.git cd intelliperf ``` diff --git a/apptainer/build.sh b/apptainer/build.sh index 6fc89cca..7954f962 100755 --- a/apptainer/build.sh +++ b/apptainer/build.sh @@ -1,69 +1,28 @@ #!/bin/bash -################################################################################ -# MIT License - -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. - -# 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: - -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. - -# 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. -################################################################################ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. debug=0 +output_path="apptainer/intelliperf.sif" while [[ $# -gt 0 ]]; do case $1 in - -d|--debug) - debug=1 - shift + -o|--output) + output_path="$2" + shift 2 ;; *) - echo "Usage: $0 [-d|--debug]" + echo "Usage: $0 [-o|--output OUTPUT_PATH]" + echo " -o, --output Specify output path for the .sif file (default: apptainer/intelliperf.sif)" exit 1 ;; esac done - -# Auto config SSH agent -if [ ! -S ~/.ssh/ssh_auth_sock ]; then - eval `ssh-agent` > /dev/null - ln -sf "$SSH_AUTH_SOCK" ~/.ssh/ssh_auth_sock -fi -export SSH_AUTH_SOCK=~/.ssh/ssh_auth_sock -[ -f ~/.ssh/id_rsa ] && ssh-add ~/.ssh/id_rsa -[ -f ~/.ssh/id_ed25519 ] && ssh-add ~/.ssh/id_ed25519 - -ssh_auth_sock_path=$(readlink -f "$SSH_AUTH_SOCK") -# Build the Singularity container -# --build-arg SSH_AUTH_SOCK=$SSH_AUTH_SOCK is used to pass the SSH agent socket to the container -# (advantage of this method is that the key is at no point copied to the container image.) -# If your SSH_AUTH_SOCK will not already bound to the container, and is available at /run/..., add `--bind /run` to the build command definition="apptainer/intelliperf.def" -if [[ $debug -eq 1 ]]; then - image="apptainer/intelliperf_debug.sif" - cmake_build_type="Debug" -else - image="apptainer/intelliperf.sif" - cmake_build_type="Release" -fi +echo "Building Apptainer image..." +echo "Definition file: $definition" +echo "Output path: $output_path" -apptainer build \ - --build-arg SSH_AUTH_SOCK=${ssh_auth_sock_path} \ - --build-arg CMAKE_BUILD_TYPE=${cmake_build_type}\ - $image $definition \ No newline at end of file +apptainer build $output_path $definition \ No newline at end of file diff --git a/apptainer/intelliperf.def b/apptainer/intelliperf.def index 28c6d4af..140a8c11 100644 --- a/apptainer/intelliperf.def +++ b/apptainer/intelliperf.def @@ -10,14 +10,6 @@ From: ubuntu:22.04 export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH export ROCM_PATH=/opt/rocm - # Misc globals - export GT_TUNING=/root/guided-tuning - export PATH=/opt/omniprobe/bin/logDuration:$PATH - export PATH=/root/rocprofiler-compute/src:$PATH - -%files - examples/bank_conflict/llm.c/requirements.txt /examples/bank_conflict/llm.c/requirements.txt - %post # Set locale apt-get -y update @@ -35,13 +27,6 @@ From: ubuntu:22.04 python3 -m pip install --upgrade pip python3 -m pip install 'cmake==3.22' - # Add GitHub trusted host - mkdir -p ~/.ssh - touch ~/.ssh/known_hosts - ssh-keyscan github.com >> ~/.ssh/known_hosts - chmod 700 ~/.ssh - chmod 644 ~/.ssh/known_hosts - # Install ROCm apt-get -y update wget https://repo.radeon.com/amdgpu-install/6.3.3/ubuntu/jammy/amdgpu-install_6.3.60303-1_all.deb @@ -52,47 +37,5 @@ From: ubuntu:22.04 export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH export ROCM_PATH=/opt/rocm - # Install rocprof-compute (via package manager) - # python3 -m pip install --ignore-installed blinker - # python3 -m pip install -r /opt/rocm/libexec/rocprofiler-compute/requirements.txt - # Install rocprof-compute (from feature branch) - export SSH_AUTH_SOCK={{ SSH_AUTH_SOCK }} - cd /root - git clone -v https://github.com/ROCm/rocprofiler-compute.git - cd rocprofiler-compute - git checkout 41e73650d5cfc3dbd98e007d6279235578f8529a - python3 -m pip install --ignore-installed blinker - python3 -m pip install -r requirements.txt - cd src - export PATH=$PWD:$PATH - - # Install Triton (version pinned) - cd /root - export TRITON_HOME=/root - git clone -v https://github.com/triton-lang/triton.git - cd triton - git checkout 6fa33ef1eecc97348d056688df84845db7d22507 - python3 -m pip install ninja wheel pybind11 - python3 -m pip install -e python - - # Install omniprobe - echo "Building with CMAKE_BUILD_TYPE={{ CMAKE_BUILD_TYPE }}" - cd /root - git clone -v git@github.com:AARInternal/omniprobe.git - cd omniprobe - git checkout 9083730ab0da50114c767773df49cb1d2165ba7f - git submodule update --init --recursive - mkdir -p build - cmake -DCMAKE_INSTALL_PREFIX=/opt/omniprobe\ - -DCMAKE_PREFIX_PATH=${ROCM_PATH}\ - -DTRITON_LLVM=/root/.triton/llvm/llvm-ubuntu-x64\ - -DCMAKE_BUILD_TYPE={{ CMAKE_BUILD_TYPE }}\ - -DCMAKE_VERBOSE_MAKEFILE=ON -S . -B build - cmake --build build --target install - export PATH=/opt/omniprobe/bin/logDuration:$PATH - - # Install agents dependencies - python3 -m pip install openai - - # Install examples dependencies - pip3 install --no-cache-dir -r /examples/bank_conflict/llm.c/requirements.txt + # Install PyTorch ROCm + python3 -m pip install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.3 \ No newline at end of file diff --git a/apptainer/run.sh b/apptainer/run.sh index 9eb9f0d7..eafca83c 100755 --- a/apptainer/run.sh +++ b/apptainer/run.sh @@ -1,27 +1,6 @@ #!/bin/bash -################################################################################ -# MIT License - -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. - -# 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: - -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. - -# 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. -################################################################################ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" parent_dir="$(dirname "$script_dir")" @@ -38,12 +17,8 @@ while [[ $# -gt 0 ]]; do size=$2 shift 2 ;; - -d|--debug) - debug=1 - shift - ;; *) - echo "Usage: $0 [-s size] [-d|--debug]" + echo "Usage: $0 [-s size]" exit 1 ;; esac @@ -59,9 +34,5 @@ fi echo "[Log] Utilize the directory /var/cache/intelliperf as a sandbox to store data you'd like to persist between container runs." # Run the container -if [[ $debug -eq 1 ]]; then - image="apptainer/intelliperf_debug.sif" -else - image="apptainer/intelliperf.sif" -fi -apptainer exec --bind $HOME/.ssh:/root/.ssh:ro --overlay ${overlay} --pwd "$working_dir" --cleanenv --env OPENAI_API_KEY=$OPENAI_API_KEY $image bash --rcfile /etc/bash.bashrc +image="apptainer/intelliperf.sif" +apptainer exec --overlay ${overlay} --pwd "$working_dir" --cleanenv --env LLM_GATEWAY_KEY=$LLM_GATEWAY_KEY $image bash --rcfile /etc/bash.bashrc diff --git a/apptainer/run_cmd.sh b/apptainer/run_cmd.sh index 928eab38..f852c4b9 100755 --- a/apptainer/run_cmd.sh +++ b/apptainer/run_cmd.sh @@ -1,27 +1,6 @@ #!/bin/bash -################################################################################ -# MIT License - -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. - -# 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: - -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. - -# 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. -################################################################################ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" parent_dir="$(dirname "$script_dir")" diff --git a/docker/build.sh b/docker/build.sh index c91ac2cd..5a445ae1 100755 --- a/docker/build.sh +++ b/docker/build.sh @@ -1,42 +1,6 @@ #!/bin/bash -################################################################################ -# MIT License - -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. - -# 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: - -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. - -# 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. -################################################################################ - -# Parse command line arguments -dev_mode=false -while [[ $# -gt 0 ]]; do - case $1 in - --dev|-d) - dev_mode=true - shift - ;; - *) - echo "Unknown option: $1" - exit 1 - ;; - esac -done +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. # Container name name="intelliperf" @@ -45,36 +9,11 @@ name="intelliperf" script_dir="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" parent_dir="$(dirname "$script_dir")" cur_dir=$(pwd) - -# Set INTELLIPERF_HOME based on dev mode -if [ "$dev_mode" = true ]; then - intelliperf_home="$cur_dir" -else - intelliperf_home="/intelliperf" -fi - pushd "$script_dir" -# Auto-configure SSH agent -if [ ! -S ~/.ssh/ssh_auth_sock ]; then - eval "$(ssh-agent)" > /dev/null - ln -sf "$SSH_AUTH_SOCK" ~/.ssh/ssh_auth_sock -fi -export SSH_AUTH_SOCK=~/.ssh/ssh_auth_sock - -# Add default keys if they exist -[ -f ~/.ssh/id_rsa ] && ssh-add ~/.ssh/id_rsa -[ -f ~/.ssh/id_ed25519 ] && ssh-add ~/.ssh/id_ed25519 -[ -f ~/.ssh/id_github ] && ssh-add ~/.ssh/id_github - -# Enable BuildKit and build the Docker image -export DOCKER_BUILDKIT=1 docker build \ - --ssh default \ -t "$name" \ - --build-arg DEV_MODE="$dev_mode" \ - --build-arg INTELLIPERF_HOME="$intelliperf_home" \ -f "$script_dir/intelliperf.Dockerfile" \ . -popd +popd \ No newline at end of file diff --git a/docker/intelliperf.Dockerfile b/docker/intelliperf.Dockerfile index b7a3adc2..28b87ea4 100644 --- a/docker/intelliperf.Dockerfile +++ b/docker/intelliperf.Dockerfile @@ -1,8 +1,5 @@ -# syntax=docker/dockerfile:1.4 - FROM rocm/vllm-dev:nightly_aiter_integration_final_20250325 -ARG DEV_MODE=false ARG INTELLIPERF_HOME=/intelliperf ENV LANG=en_US.UTF-8 @@ -19,19 +16,5 @@ RUN apt-get update && apt-get install -y \ gdb \ && locale-gen en_US.UTF-8 -# Add GitHub trusted host -RUN mkdir -p ~/.ssh && \ - touch ~/.ssh/known_hosts && \ - ssh-keyscan github.com >> ~/.ssh/known_hosts && \ - chmod 700 ~/.ssh && \ - chmod 644 ~/.ssh/known_hosts - # Set the working directory -WORKDIR $INTELLIPERF_HOME - -# Clone IntelliPerf only in non-dev mode -RUN --mount=type=ssh bash -c 'if [ "$DEV_MODE" = "false" ]; then \ - git clone git@github.com:AMDResearch/intelliperf.git . ; \ - pip install -e .; \ - python3 scripts/install_tool.py --all; \ - fi' +WORKDIR $INTELLIPERF_HOME \ No newline at end of file diff --git a/docker/run.sh b/docker/run.sh index 4f4ecdae..118d0a6b 100755 --- a/docker/run.sh +++ b/docker/run.sh @@ -1,27 +1,6 @@ #!/bin/bash -################################################################################ -# MIT License - -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. - -# 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: - -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. - -# 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. -################################################################################ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. name="intelliperf" @@ -30,11 +9,7 @@ docker run -it --rm \ --device=/dev/kfd \ --device=/dev/dri \ --group-add video \ - -v $HOME/.ssh:/tmp/ssh:ro \ -v $(pwd):$(pwd) \ -w $(pwd) \ -e LLM_GATEWAY_KEY="$LLM_GATEWAY_KEY" \ - -e SSH_AUTH_SOCK="$SSH_AUTH_SOCK" \ - -v $SSH_AUTH_SOCK:$SSH_AUTH_SOCK \ - "$name" \ - bash -c "cp -r /tmp/ssh/* /root/.ssh/ 2>/dev/null || true && chown -R root:root /root/.ssh && chmod 700 /root/.ssh && chmod 600 /root/.ssh/config /root/.ssh/id_* /root/.ssh/known_hosts 2>/dev/null || true; exec bash" + "$name" \ No newline at end of file diff --git a/pyproject.toml b/pyproject.toml index 6a2bdc8c..36d735f3 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -12,7 +12,7 @@ readme = "README.md" requires-python = ">=3.8" # Python dependencies -dependencies = ["tomli", "openai>=1.0.0", "tabulate", "ml_dtypes", "dspy", "pandas", "duckdb", "rich", "pytest"] +dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==3.0.3", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]"] [tool.setuptools] @@ -31,7 +31,7 @@ build-backend = "setuptools.build_meta" [tool.rocprofiler-compute] git = "https://github.com/ROCm/rocprofiler-compute" -branch = "41e73650d5cfc3dbd98e007d6279235578f8529a" +branch = "a3dc98e25106f0cfb5f996acc9848274b4e6cf15" build_command = """ python3 -m pip install --ignore-installed blinker && python3 -m pip install -r requirements.txt From 456b3ea5ef75b4ef19c5f3da4c52a560adafcf35 Mon Sep 17 00:00:00 2001 From: JoseSantosAMD <87447437+JoseSantosAMD@users.noreply.github.com> Date: Tue, 30 Sep 2025 18:48:56 -0400 Subject: [PATCH 03/14] Add rpds-py dependency (#147) --- pyproject.toml | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/pyproject.toml b/pyproject.toml index 36d735f3..468492ec 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -12,7 +12,7 @@ readme = "README.md" requires-python = ">=3.8" # Python dependencies -dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==3.0.3", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]"] +dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==3.0.3", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]", "rpds-py"] [tool.setuptools] From 06cca45fd28f34cda28783842c82c10a7e92f8a0 Mon Sep 17 00:00:00 2001 From: Arya Tschand Date: Tue, 7 Oct 2025 18:17:14 -0400 Subject: [PATCH 04/14] SwizzlePerf (#143) Co-authored-by: github-actions[bot] Co-authored-by: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Co-authored-by: Muhammad Awad Co-authored-by: stephen youn <13525892+stephen-youn@users.noreply.github.com> --- .github/workflows/ci.yml | 2 +- .../workflows/scripts/check_test_results.sh | 1 + .github/workflows/scripts/ci_tests.sh | 2 +- .gitignore | 1 + apptainer/build.sh | 3 +- apptainer/intelliperf.def | 66 ++- examples/triton/gemm.py | 205 +++++++ examples/triton/gemm_runner.py | 226 ++++++++ pyproject.toml | 5 +- src/intelliperf/__main__.py | 47 +- src/intelliperf/core/gpu_spec.py | 53 ++ src/intelliperf/core/llm.py | 123 +++- src/intelliperf/formulas/atomic_contention.py | 3 +- src/intelliperf/formulas/bank_conflict.py | 3 +- src/intelliperf/formulas/diagnose_only.py | 2 +- src/intelliperf/formulas/formula_base.py | 20 +- src/intelliperf/formulas/memory_access.py | 3 +- src/intelliperf/formulas/swizzling.py | 532 ++++++++++++++++++ 18 files changed, 1219 insertions(+), 78 deletions(-) create mode 100755 examples/triton/gemm.py create mode 100755 examples/triton/gemm_runner.py create mode 100644 src/intelliperf/formulas/swizzling.py diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index b6a24fbe..6eb59014 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -43,7 +43,7 @@ jobs: name: IntelliPerf Test needs: build-apptainer-image runs-on: [self-hosted] - timeout-minutes: 20 + timeout-minutes: 60 steps: - name: Checkout repository diff --git a/.github/workflows/scripts/check_test_results.sh b/.github/workflows/scripts/check_test_results.sh index cb91e0e5..00821a0a 100755 --- a/.github/workflows/scripts/check_test_results.sh +++ b/.github/workflows/scripts/check_test_results.sh @@ -29,6 +29,7 @@ overall_success=true check_test_result "$results_dir/memory_access_output.json" "Memory Access" || overall_success=false check_test_result "$results_dir/bank_conflict_output.json" "Bank Conflict" || overall_success=false check_test_result "$results_dir/atomic_contention_output.json" "Atomic Contention" || overall_success=false +check_test_result "$results_dir/swizzling_output.json" "Swizzling" || overall_success=false check_test_result "$results_dir/diagnose_only_hip_uncoalesced.json" "Diagnose Only (HIP)" || overall_success=false check_test_result "$results_dir/diagnose_only_torch_add.json" "Diagnose Only (Torch)" || overall_success=false check_test_result "$results_dir/diagnose_only_triton_reduce.json" "Diagnose Only (Triton)" || overall_success=false diff --git a/.github/workflows/scripts/ci_tests.sh b/.github/workflows/scripts/ci_tests.sh index 4916838b..34848ea7 100755 --- a/.github/workflows/scripts/ci_tests.sh +++ b/.github/workflows/scripts/ci_tests.sh @@ -16,7 +16,7 @@ model="openai/gpt-4o" intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --formula=memoryAccess -o $results_dir/memory_access_output.json -- ./build/access_pattern/uncoalesced || true intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --formula=bankConflict -o $results_dir/bank_conflict_output.json -- ./build/bank_conflict/matrix_transpose 1024 1024 || true intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --instrument_command="./scripts/build_examples.sh -i -c" --formula=atomicContention -o $results_dir/atomic_contention_output.json -- ./build/contention/reduction || true - +intelliperf -vvv --project_directory=./examples --provider $provider --model $model --formula=swizzling --project_directory="./examples" --unittest_command="triton/gemm_runner.py --validate" -o $results_dir/swizzling_output.json -- ./triton/gemm_runner.py || true # Diagnose Only intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_hip_uncoalesced.json -- ./examples/build/access_pattern/uncoalesced intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_torch_add.json -- ./examples/torch/add.py diff --git a/.gitignore b/.gitignore index c3a13cb1..41639bab 100644 --- a/.gitignore +++ b/.gitignore @@ -13,3 +13,4 @@ external/ intelliperf_env/ trace/ +.build/ diff --git a/apptainer/build.sh b/apptainer/build.sh index 7954f962..847006e5 100755 --- a/apptainer/build.sh +++ b/apptainer/build.sh @@ -25,4 +25,5 @@ echo "Building Apptainer image..." echo "Definition file: $definition" echo "Output path: $output_path" -apptainer build $output_path $definition \ No newline at end of file +# Build the container +apptainer build --force $output_path $definition \ No newline at end of file diff --git a/apptainer/intelliperf.def b/apptainer/intelliperf.def index 140a8c11..da1fa972 100644 --- a/apptainer/intelliperf.def +++ b/apptainer/intelliperf.def @@ -1,41 +1,65 @@ Bootstrap: docker -From: ubuntu:22.04 +From: rocm/pytorch:rocm7.0_ubuntu22.04_py3.10_pytorch_release_2.8.0 %environment # Locale export LANG=en_US.UTF-8 # ROCm globals - export PATH=/opt/rocm/bin:$PATH + export PATH=/opt/conda/envs/py_3.10/bin:/opt/rocm/bin:$PATH export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH export ROCM_PATH=/opt/rocm + export OMPI_MCA_mtl="^ofi" + export OMPI_MCA_pml="ob1" %post + /bin/bash -c " # Set locale apt-get -y update apt-get install -y locales locale-gen en_US.UTF-8 export LANG=en_US.UTF-8 - # Install dependencies + # Install additional system dependencies apt-get -y update - apt-get install -y software-properties-common - apt-get upgrade -y - apt-get install -y build-essential python3 python3-pip python3-setuptools python3-wheel git wget clang lld libzstd-dev libomp-dev vim libdwarf-dev - apt-get install -y locales - locale-gen en_US.UTF-8 - python3 -m pip install --upgrade pip - python3 -m pip install 'cmake==3.22' + apt-get install -y software-properties-common git wget clang lld libzstd-dev libomp-dev vim libdwarf-dev gdb tmux - # Install ROCm - apt-get -y update - wget https://repo.radeon.com/amdgpu-install/6.3.3/ubuntu/jammy/amdgpu-install_6.3.60303-1_all.deb - apt-get -y install ./amdgpu-install_6.3.60303-1_all.deb - apt-get -y update - apt-get install -y rocm-dev rocm-llvm-dev rocm-hip-runtime-dev rocm-smi-lib rocminfo rocthrust-dev rocprofiler-compute rocblas rocm-gdb gdb tmux - export PATH=/opt/rocm/bin:$PATH - export LD_LIBRARY_PATH=/opt/rocm/lib:$LD_LIBRARY_PATH - export ROCM_PATH=/opt/rocm + # Upgrade pip + pip install --upgrade pip + + # Install additional Python packages for intelliperf + pip install --no-cache-dir \ + astunparse==1.6.2 \ + colorlover \ + dash-bootstrap-components \ + dash-svg \ + dash>=3.0.0 \ + kaleido==0.2.1 \ + matplotlib \ + numpy>=1.17.5 \ + pandas>=1.4.3 \ + plotext \ + plotille \ + pymongo \ + pyyaml \ + setuptools \ + tabulate \ + textual \ + textual_plotext \ + textual-fspicker \ + tqdm \ + tomli\ + tabulate\ + ml_dtypes\ + dspy==2.6.27\ + pandas\ + duckdb\ + rich\ + pytest\ + litellm[proxy] + " - # Install PyTorch ROCm - python3 -m pip install --pre torch torchvision torchaudio --index-url https://download.pytorch.org/whl/nightly/rocm6.3 \ No newline at end of file +%runscript + echo "Welcome to IntelliPerf with ROCm 7.0!" + source /opt/conda/bin/activate py_3.10 + exec "$@" diff --git a/examples/triton/gemm.py b/examples/triton/gemm.py new file mode 100755 index 00000000..9467ab5e --- /dev/null +++ b/examples/triton/gemm.py @@ -0,0 +1,205 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +import triton +import triton.language as tl + + +@triton.jit() +def streamk_gemm( + A, + B, + C, + bias_ptr, + P, + locks, + M, + N, + K, + stride_am, + stride_ak, + stride_bk, + stride_bn, + stride_cm, + stride_cn, + stride_bias, + BLOCK_SIZE_M: tl.constexpr, + BLOCK_SIZE_N: tl.constexpr, + BLOCK_SIZE_K: tl.constexpr, + GROUP_SIZE_M: tl.constexpr, + NUM_SMS: tl.constexpr, + STREAMK_TILES: tl.constexpr, + NUM_XCDS: tl.constexpr, + BIAS: tl.constexpr, + EVEN_K: tl.constexpr, +): + # original program‐ID + pid = tl.program_id(0) + + num_pid_m = tl.cdiv(M, BLOCK_SIZE_M) + num_pid_n = tl.cdiv(N, BLOCK_SIZE_N) + iters_per_tile = tl.cdiv(K, BLOCK_SIZE_K) + total_tiles = num_pid_m * num_pid_n + total_full_tiles = total_tiles - STREAMK_TILES + + tl.assume(stride_am > 0) + tl.assume(stride_ak > 0) + tl.assume(stride_bn > 0) + tl.assume(stride_bk > 0) + tl.assume(stride_cm > 0) + tl.assume(stride_cn > 0) + + acc_dtype = tl.float32 if C.type.element_ty != tl.int8 else tl.int32 + + for tile_id in range(pid, total_full_tiles, NUM_SMS): + num_pid_in_group = GROUP_SIZE_M * num_pid_n + group_id = tile_id // num_pid_in_group + first_pid_m = group_id * GROUP_SIZE_M + group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M) + pid_in_group = tile_id % num_pid_in_group + + pid_m = first_pid_m + pid_in_group % group_size_m + pid_n = pid_in_group // group_size_m + tl.assume(pid_m >= 0) + tl.assume(pid_n >= 0) + + rm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + rn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) + rm = rm % M + rn = rn % N + + rk = tl.arange(0, BLOCK_SIZE_K) + A_BASE = A + rm[:, None] * stride_am + rk[None, :] * stride_ak + B_BASE = B + rk[:, None] * stride_bk + rn[None, :] * stride_bn + + if BIAS: + bias_ = bias_ptr + rm * stride_bias + bias = tl.load(bias_, mask=rm < M, other=0.0) + + loop_k = tl.cdiv(K, BLOCK_SIZE_K) + if not EVEN_K: + loop_k -= 1 + + acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=acc_dtype) + for k in range(0, loop_k): + a = tl.load(tl.multiple_of(A_BASE, (1, 16))) + b = tl.load(tl.multiple_of(B_BASE, (16, 1))) + acc += tl.dot(a, b) + A_BASE += BLOCK_SIZE_K * stride_ak + B_BASE += BLOCK_SIZE_K * stride_bk + + if not EVEN_K: + k = loop_k + rk = k * BLOCK_SIZE_K + tl.arange(0, BLOCK_SIZE_K) + A_BASE = A + rm[:, None] * stride_am + rk[None, :] * stride_ak + B_BASE = B + rk[:, None] * stride_bk + rn[None, :] * stride_bn + A_BASE = tl.multiple_of(A_BASE, (1, 16)) + B_BASE = tl.multiple_of(B_BASE, (16, 1)) + a = tl.load(A_BASE, mask=rk[None, :] < K, other=0.0) + b = tl.load(B_BASE, mask=rk[:, None] < K, other=0.0) + acc += tl.dot(a, b) + + c = acc.to(C.type.element_ty) + if BIAS: + c += bias[:, None] + + pid_m = first_pid_m + pid_in_group % group_size_m + pid_n = pid_in_group // group_size_m + rm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + rn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) + rm = rm % M + rn = rn % N + C_ = C + rm[:, None] * stride_cm + rn[None, :] * stride_cn + mask = (rm < M)[:, None] & (rn < N)[None, :] + tl.store(C_, c, mask=mask) + + tl.assume(pid >= 0) + total_streamk_iters = STREAMK_TILES * iters_per_tile + streamk_iters_pcu = total_streamk_iters // NUM_SMS + streamk_remainder_iters = total_streamk_iters % NUM_SMS + start_iter = total_full_tiles * iters_per_tile + pid * streamk_iters_pcu + tl.minimum(pid, streamk_remainder_iters) + last_iter = ( + total_full_tiles * iters_per_tile + (pid + 1) * streamk_iters_pcu + tl.minimum(pid + 1, streamk_remainder_iters) + ) + while start_iter < last_iter: + remainder = start_iter % iters_per_tile + end_iter = tl.minimum(start_iter + (iters_per_tile - remainder), last_iter) + tile_id = start_iter // iters_per_tile + num_pid_in_group = GROUP_SIZE_M * num_pid_n + group_id = tile_id // num_pid_in_group + first_pid_m = group_id * GROUP_SIZE_M + group_size_m = min(num_pid_m - first_pid_m, GROUP_SIZE_M) + pid_in_group = tile_id % num_pid_in_group + + pid_m = first_pid_m + pid_in_group % group_size_m + pid_n = pid_in_group // group_size_m + tl.assume(pid_m >= 0) + tl.assume(pid_n >= 0) + + rm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + rn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) + rm = rm % M + rn = rn % N + rk = tl.arange(0, BLOCK_SIZE_K) + A_BASE = A + rm[:, None] * stride_am + rk[None, :] * stride_ak + BLOCK_SIZE_K * stride_ak * remainder + B_BASE = B + rk[:, None] * stride_bk + rn[None, :] * stride_bn + BLOCK_SIZE_K * stride_bk * remainder + A_BASE = tl.multiple_of(A_BASE, (1, 16)) + B_BASE = tl.multiple_of(B_BASE, (16, 1)) + + if BIAS: + bias_ = bias_ptr + rm * stride_bias + bias = tl.load(bias_, mask=rm < M, other=0.0) + + acc = tl.zeros((BLOCK_SIZE_M, BLOCK_SIZE_N), dtype=acc_dtype) + for current_iter in range(start_iter, end_iter): + if EVEN_K: + a = tl.load(A_BASE) + b = tl.load(B_BASE) + else: + global_k_offset = (current_iter % iters_per_tile) * BLOCK_SIZE_K + k_mask = global_k_offset + rk < K + a = tl.load(A_BASE, mask=k_mask[None, :], other=0.0) + b = tl.load(B_BASE, mask=k_mask[:, None], other=0.0) + acc += tl.dot(a, b) + A_BASE += BLOCK_SIZE_K * stride_ak + B_BASE += BLOCK_SIZE_K * stride_bk + + tile_iter = tile_id * iters_per_tile + + if start_iter != tile_iter: + rm1 = tl.arange(0, BLOCK_SIZE_M) + rn1 = tl.arange(0, BLOCK_SIZE_N) + P_ = P + pid * BLOCK_SIZE_M * BLOCK_SIZE_N + rm1[:, None] * BLOCK_SIZE_N + rn1[None, :] + tl.store(P_, acc, cache_modifier=".wt") + tl.debug_barrier() + tl.store(locks + pid, 1, cache_modifier=".wt") + else: + next_pid = pid + 1 + tile_iter_end = tile_iter + iters_per_tile + end = end_iter + while end < tile_iter_end and next_pid < NUM_SMS: + while tl.load(locks + next_pid, cache_modifier=".cv", volatile=True) != 1: + pass + rm1 = tl.arange(0, BLOCK_SIZE_M) + rn1 = tl.arange(0, BLOCK_SIZE_N) + P_ = P + next_pid * BLOCK_SIZE_M * BLOCK_SIZE_N + rm1[:, None] * BLOCK_SIZE_N + rn1[None, :] + acc += tl.load(tl.multiple_of(P_, (1, 16)), cache_modifier=".cv") + end += streamk_iters_pcu + (next_pid < streamk_remainder_iters) + next_pid += 1 + + c = acc.to(C.type.element_ty) + if BIAS: + c += bias[:, None] + + pid_m = first_pid_m + pid_in_group % group_size_m + pid_n = pid_in_group // group_size_m + rm = pid_m * BLOCK_SIZE_M + tl.arange(0, BLOCK_SIZE_M) + rn = pid_n * BLOCK_SIZE_N + tl.arange(0, BLOCK_SIZE_N) + rm = rm % M + rn = rn % N + C_ = C + rm[:, None] * stride_cm + rn[None, :] * stride_cn + mask = (rm < M)[:, None] & (rn < N)[None, :] + tl.store(C_, c, mask=mask) + + start_iter = end_iter diff --git a/examples/triton/gemm_runner.py b/examples/triton/gemm_runner.py new file mode 100755 index 00000000..f2524b35 --- /dev/null +++ b/examples/triton/gemm_runner.py @@ -0,0 +1,226 @@ +#!/usr/bin/env python3 +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +import argparse +import os +import random +import sys + +import torch +import triton + +# Ensure we can import the sibling gemm module when run as ./triton/gemm_runner.py +sys.path.append(os.path.dirname(__file__)) +from gemm import streamk_gemm # noqa: E402 + + +class matmul(torch.autograd.Function): + _debug = True + + @staticmethod + def set_debug(debug: bool): + matmul._debug = debug + + @staticmethod + def _call( + a: torch.Tensor, + b: torch.Tensor, + c: torch.Tensor, + bias: torch.Tensor, + P: torch.Tensor, + locks: torch.Tensor, + total_programs_streamk: int, + BLK_M: int, + BLK_N: int, + BLK_K: int, + gsize_m: int, + two_tiles: bool, + num_stages: int, + num_warps: int, + waves_per_eu: int, + mfmaInstrSize: int, + kpack: int, + ): + assert a.shape[1] == b.shape[0], "incompatible dimensions" + M, K = a.shape + _, N = b.shape + + total_blocks_M = triton.cdiv(M, BLK_M) + total_blocks_N = triton.cdiv(N, BLK_N) + iters_per_tile = triton.cdiv(K, BLK_K) + total_tiles = total_blocks_M * total_blocks_N + even_k = K % BLK_K == 0 + + if total_programs_streamk > 0: + total_tiles_streamk = total_tiles % total_programs_streamk + total_blocking_tiles = total_tiles - total_tiles_streamk + total_iters_streamk = total_tiles_streamk * iters_per_tile + total_full_tiles_streamk = total_iters_streamk // total_programs_streamk + total_partial_tiles_streamk = total_iters_streamk % total_programs_streamk + else: + total_blocking_tiles = total_tiles + total_tiles_streamk = 0 + total_full_tiles_streamk = 0 + total_partial_tiles_streamk = 0 + total_iters_streamk = 0 + + if matmul._debug: + print(f"M,N,K={M},{N},{K} ; BLK_M,N,K={BLK_M},{BLK_N},{BLK_K}") + print(f"{total_blocks_M=} x {total_blocks_N=} = {total_tiles=}") + print(f"{total_tiles_streamk=} + {total_blocking_tiles=} = {total_tiles=}") + print(f"{total_programs_streamk=}") + print(f"{total_blocking_tiles=}") + print(f"{total_full_tiles_streamk=}") + print(f"{iters_per_tile=}") + print(f"{total_iters_streamk=}") + print("total_remainder_iters_streamk=", total_partial_tiles_streamk) + + use_bias = False + grids = total_programs_streamk + stride_bias = bias.stride(0) if use_bias else 0 + num_xcds = 8 + + kk = streamk_gemm[(grids,)]( + a, + b, + c, + bias, + P, + locks, + M, + N, + K, + a.stride(0), + a.stride(1), + b.stride(0), + b.stride(1), + c.stride(0), + c.stride(1), + stride_bias, + BLOCK_SIZE_M=BLK_M, + BLOCK_SIZE_N=BLK_N, + BLOCK_SIZE_K=BLK_K, + GROUP_SIZE_M=gsize_m, + NUM_SMS=total_programs_streamk, + STREAMK_TILES=total_tiles_streamk, + NUM_XCDS=num_xcds, + BIAS=use_bias, + EVEN_K=even_k, + ) + if matmul._debug: + print(f"{kk.n_regs} registers used, {kk.n_spills} spills") + + return c + + @staticmethod + def forward( + ctx, + a: torch.Tensor, + b: torch.Tensor, + c: torch.Tensor, + bias: torch.Tensor, + P: torch.Tensor, + locks: torch.Tensor, + grid: int, + BLK_M=128, + BLK_N=128, + BLK_K=32, + gsize_m=1, + two_tiles=True, + num_stages=3, + num_warps=4, + waves_per_eu=2, + mfmaInstrSize=16, + kpack=1, + ): + matmul._call( + a=a, + b=b, + c=c, + bias=bias, + P=P, + locks=locks, + total_programs_streamk=grid, + BLK_M=BLK_M, + BLK_N=BLK_N, + BLK_K=BLK_K, + gsize_m=gsize_m, + two_tiles=two_tiles, + num_warps=num_warps, + num_stages=num_stages, + waves_per_eu=waves_per_eu, + mfmaInstrSize=mfmaInstrSize, + kpack=kpack, + ) + return c + + +if __name__ == "__main__": + parser = argparse.ArgumentParser() + parser.add_argument("--validate", action="store_true", help="Validate the Triton implementation against PyTorch.") + args = parser.parse_args() + + torch.manual_seed(123) + random.seed(123) + + total_sm = 304 + print(f"total SMs: {total_sm}") + + m, n, k = 8192, 8192, 8192 + A = torch.randn(m, k, device="cuda", dtype=torch.float16) + B = torch.randn(n, k, device="cuda", dtype=torch.float16).T + C = torch.zeros((m, n), device="cuda", dtype=A.dtype) + bias = torch.zeros((m,), device="cuda", dtype=A.dtype) + + BLK_M = 256 + BLK_N = 256 + BLK_K = 64 + + total_blocks_M = triton.cdiv(m, BLK_M) + total_blocks_N = triton.cdiv(n, BLK_N) + total_tiles = total_blocks_M * total_blocks_N + + gsize_m = 8 + two_tiles = "True" + num_stages = 2 + num_warps = 8 + waves_per_eu = 0 + mfmaInstrSize = 16 + kpack = 2 + + print(f"{total_sm=}") + matmul.set_debug(True) + locks = torch.zeros((total_sm,), device="cuda", dtype=torch.int32) + P = torch.zeros((total_sm, BLK_M * BLK_N), device="cuda", dtype=torch.float32) + + C = matmul.apply( + A, + B, + C, + bias, + P, + locks, + total_sm, + BLK_M, + BLK_N, + BLK_K, + gsize_m, + two_tiles, + num_stages, + num_warps, + waves_per_eu, + mfmaInstrSize, + kpack, + ) + + if args.validate: + expected = A @ B + if not torch.allclose(C, expected, atol=1): + max_diff = (C - expected).abs().max().item() + print(f"Validation Failed: max abs diff = {max_diff}") + sys.exit(1) + else: + print("Validation Successful!") + else: + print("Completed GEMM run.") diff --git a/pyproject.toml b/pyproject.toml index 468492ec..dc375f01 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -9,11 +9,10 @@ authors = [ license = { text = "MIT" } readme = "README.md" -requires-python = ">=3.8" +requires-python = ">=3.10" # Python dependencies -dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==3.0.3", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]", "rpds-py"] - +dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==2.6.27", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]", "rpds-py"] [tool.setuptools] package-dir = {"" = "src"} diff --git a/src/intelliperf/__main__.py b/src/intelliperf/__main__.py index 14d7432d..bb731769 100644 --- a/src/intelliperf/__main__.py +++ b/src/intelliperf/__main__.py @@ -88,11 +88,17 @@ def intelliperf_parser(): optional_args.add_argument( "-f", "--formula", - choices=["bankConflict", "memoryAccess", "atomicContention", "diagnoseOnly"], + choices=[ + "bankConflict", + "memoryAccess", + "atomicContention", + "diagnoseOnly", + "swizzling", + ], default="diagnoseOnly", metavar="", type=str, - help="Specify the formula to use for optimization.\nAvailable options: bankConflict, memoryAccess, atomicContention, diagnoseOnly (default: diagnoseOnly)", + help="Specify the formula to use for optimization.\nAvailable options: bankConflict, memoryAccess, atomicContention, diagnoseOnly, swizzling (default: diagnoseOnly)", ) optional_args.add_argument( "--top_n", @@ -160,6 +166,9 @@ def intelliperf_parser(): # Output arguments optional_args.add_argument("-o", "--output_file", type=str, metavar="", help="Path to the output file") + # Output arguments + optional_args.add_argument("-k", "--kernel", type=str, metavar="", help="Kernel name to optimize") + args = parser.parse_args() # Handle internal LLM option @@ -221,24 +230,30 @@ def main(): from intelliperf.formulas.atomic_contention import atomic_contention formula = atomic_contention + elif args.formula == "swizzling": + from intelliperf.formulas.swizzling import swizzling + + formula = swizzling else: logging.error(f"Invalid formula specified. {args.formula} is not supported.") import sys sys.exit(1) - optimizer = formula( - name=generated_name, - build_command=args.build_command, - instrument_command=args.instrument_command, - project_directory=args.project_directory, - app_cmd=args.remaining, - top_n=args.top_n, - model=args.model, - provider=args.provider, - in_place=args.in_place, - unittest_command=args.unittest_command, - ) + optimizer_args = { + "name": generated_name, + "build_command": args.build_command, + "instrument_command": args.instrument_command, + "project_directory": args.project_directory, + "app_cmd": args.remaining, + "top_n": args.top_n, + "model": args.model, + "provider": args.provider, + "in_place": args.in_place, + "unittest_command": args.unittest_command, + } + + optimizer = formula(**optimizer_args) # Helper function to flush logs if tracing is enabled def flush_logs_if_enabled(): @@ -269,10 +284,10 @@ def flush_logs_if_enabled(): logging.info(f"Executing pass {attempt + 1} of {num_attempts}.") # Optimize the application based on insights from instrumentation. - optimize_result = optimizer.optimize_pass() + optimize_result = optimizer.optimize_pass(target_kernel=args.kernel) if not optimize_result: optimize_result.report_out() - logging.warning(f"Optimization pass {attempt + 1} failed. Retrying...") + logging.warning(f"{args.formula} optimization pass {attempt + 1} failed. Retrying...") flush_logs_if_enabled() # Flush after failed optimization continue diff --git a/src/intelliperf/core/gpu_spec.py b/src/intelliperf/core/gpu_spec.py index 011cb28d..20fedb72 100644 --- a/src/intelliperf/core/gpu_spec.py +++ b/src/intelliperf/core/gpu_spec.py @@ -190,6 +190,24 @@ def get_cu_count(device_id=None): return cu_count.value +def get_xcd_count(device_id=None): + """ + Return number of XCDs. Currently hardcoded. + """ + return 8 + + +def get_hip_runtime_scheduling_info(): + return ( + "HIP runtime scheduling of blocks:\n\n" + 'By default, the hardware scheduler assigns each incoming block, in order, to XCDs in a cyclic ("round-robin") sequence:\n\n' + "// pseudocode for default mapping for each block in [0, num_blocks):\n\n" + "assigned_xcd = block % number of XCDs; // execute block id on assigned XCD\n\n" + 'Once it reaches total XCD number - 1, it "wraps around" and continues assigning the next blocks to XCD 0, then XCD 1, and so on.\n\n' + 'If there are more blocks than XCDs, the scheduler effectively makes multiple "rounds," each of size number of XCDs.\n\n' + ) + + def get_wall_clock_rate(device_id=None): if device_id is None: device_id = get_device() @@ -264,6 +282,25 @@ def get_l2_cache_size_kb(device_id=None): return l2.value / 1024.0 +def get_l2_cache_size_per_xcd_kb(device_id=None): + """Return L2 cache size per XCD in KB.""" + total_l2 = get_l2_cache_size_kb(device_id) + num_xcds = get_xcd_count(device_id) + if num_xcds == 0: + return 0 + return total_l2 + + +def get_cus_per_xcd(device_id=None): + """Return number of CUs per XCD.""" + total_cus = get_cu_count(device_id) + num_xcds = get_xcd_count(device_id) + if num_xcds == 0: + return 0 + # Use integer division as CUs can't be fractional + return total_cus // num_xcds + + def measure_atomic_latency_ns(device_id=None, iters=1000000, block_size=256): return 1000 @@ -304,10 +341,26 @@ def get_l2_cache_size(self): """Return L2 cache size in KB.""" return get_l2_cache_size_kb(self.device_id) + def get_l2_cache_size_per_xcd(self): + """Return L2 cache size per XCD in KB.""" + return get_l2_cache_size_per_xcd_kb(self.device_id) + def get_num_cus(self): """Return number of compute units (multiprocessors).""" return get_cu_count(self.device_id) + def get_num_xcds(self): + """Return number of XCDs.""" + return get_xcd_count(self.device_id) + + def get_num_cus_per_xcd(self): + """Return number of CUs per XCD.""" + return get_cus_per_xcd(self.device_id) + + def get_hip_runtime_scheduling_info(self): + """Return HIP runtime scheduling info.""" + return get_hip_runtime_scheduling_info() + def get_atomic_latency(self): """Return average atomic-add latency in nanoseconds.""" return measure_atomic_latency_ns(self.device_id) diff --git a/src/intelliperf/core/llm.py b/src/intelliperf/core/llm.py index 5dc7f453..0eece144 100644 --- a/src/intelliperf/core/llm.py +++ b/src/intelliperf/core/llm.py @@ -33,6 +33,37 @@ class LLM: + def _get_model_context_length(self) -> Optional[int]: + """Query the model's max context length from the API""" + import logging + + try: + # Try to get model info from OpenRouter + if "openrouter" in self.provider.lower(): + models_url = "https://openrouter.ai/api/v1/models" + headers = {"Authorization": f"Bearer {self.api_key}"} + resp = requests.get(models_url, headers=headers, timeout=5) + if resp.status_code == 200: + models_data = resp.json().get("data", []) + for model_info in models_data: + if model_info.get("id") == self.model: + context_length = model_info.get("context_length") + if context_length: + logging.info(f"Model {self.model} max context: {context_length:,} tokens") + return context_length + + # Try to get from litellm/dspy metadata if available + if hasattr(self.lm, "model_info"): + context_length = getattr(self.lm.model_info, "max_tokens", None) + if context_length: + logging.info(f"Model {self.model} max context: {context_length:,} tokens") + return context_length + + except Exception as e: + logging.debug(f"Could not query model context length: {e}") + + return None + def __init__( self, api_key: str, @@ -53,10 +84,26 @@ def __init__( self.header = {"Ocp-Apim-Subscription-Key": api_key} else: self.use_amd = False - self.lm = dspy.LM(f"{self.provider}/{self.model}", api_key=api_key) + # Query model context and reserve ~20% for input, rest for output + max_context = self._get_model_context_length() + max_output_tokens = int(max_context * 0.8) if max_context else 4096 + # Set timeout to 10 minutes (600 seconds) + timeout_mins = 1 + self.lm = dspy.LM( + f"{self.provider}/{self.model}", + api_key=api_key, + max_tokens=max_output_tokens, + timeout=timeout_mins * 60, + ) dspy.configure(lm=self.lm) - def ask(self, user_prompt: str, record_meta: str = None) -> str: + def ask( + self, + user_prompt: str, + signature="prompt: str -> optimized_code: str", + answer_type: str = "optimized_code", + record_meta: str = None, + ): # Log the LLM interaction start if self.logger: self.logger.record( @@ -67,10 +114,11 @@ def ask(self, user_prompt: str, record_meta: str = None) -> str: "model": self.model, "provider": self.provider, "record_meta": record_meta, + "signature": str(signature), + "answer_type": answer_type, }, ) - # Initialize reasoning variable reasoning = None try: @@ -88,43 +136,64 @@ def ask(self, user_prompt: str, record_meta: str = None) -> str: resp = requests.post(url, json=body, headers=self.header) resp.raise_for_status() response_content = resp.json()["choices"][0]["message"]["content"] - else: - # DSPy path: use ChainOfThought with clear signature - # Define signature mapping input prompt to optimized code + + # Log successful response + if self.logger: + self.logger.record( + "llm_call_success", + { + "response": response_content, + "response_length": len(response_content), + "record_meta": record_meta, + }, + ) + + return response_content + else: # DSPy path dspy.context(description=self.system_prompt) - signature = "prompt: str -> optimized_code: str" chain = dspy.ChainOfThought(signature) ct_response = chain(prompt=user_prompt) - # Extract both the reasoning and the final answer - response_content = getattr(ct_response, "optimized_code", str(ct_response)) - - # Try to capture the reasoning/chain-of-thought steps + # Try to capture reasoning if available (not returned) reasoning = getattr(ct_response, "reasoning", None) - # Log successful response with reasoning if available - if self.logger: - success_data = { - "response": response_content, - "response_length": len(response_content), - "record_meta": record_meta, - } - if reasoning: - success_data["reasoning"] = reasoning - success_data["reasoning_type"] = "chain_of_thought" - - self.logger.record("llm_call_success", success_data) - - return response_content + # Determine what to return based on signature type + if isinstance(signature, str): + # Simple signature: extract the requested answer_type field + response_content = getattr(ct_response, answer_type, str(ct_response)) + else: + # Complex signature (e.g., dspy.Signature subclass): return full prediction object + response_content = ct_response + + # Log successful response + if self.logger: + log_payload = { + "record_meta": record_meta, + "signature": str(signature), + "answer_type": answer_type, + } + try: + if isinstance(response_content, str): + log_payload["response"] = response_content + log_payload["response_length"] = len(response_content) + else: + # Best-effort: record fields present on prediction object + log_payload["response_fields"] = list(getattr(ct_response, "__dict__", {}).keys()) + except Exception: + pass + if reasoning: + log_payload["reasoning"] = reasoning + log_payload["reasoning_type"] = "chain_of_thought" + self.logger.record("llm_call_success", log_payload) + + return response_content except Exception as e: error_message = str(e) error_type = type(e).__name__ - if self.logger: self.logger.record( "llm_call_error", {"error": error_message, "error_type": error_type, "record_meta": record_meta} ) - print(f"ERROR: {error_message}") sys.exit(1) diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index 57035e9e..8bc8f68a 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -130,7 +130,7 @@ def build_pass(self, validate_build_result=True) -> Result: self.current_summary = result.error_report return result - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000) -> Result: + def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: """ Optimize the kernel to remove atomic contention via OpenAI API @@ -177,6 +177,7 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000) -> Res field=field, subfield=subfield, comparison_func=lambda x: x > average_atomic_lat, + target_kernel=target_kernel, ) if len(filtered_report_card) == 0: diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index 57ff6bbe..95648491 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -201,7 +201,7 @@ def instrument_pass(self) -> Result: ) return Result(success=True, asset=self._instrumentation_results) - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000) -> Result: + def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: """ Optimize the kernel to remove shared memory bank conflicts via OpenAI API @@ -248,6 +248,7 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000) -> Res field="lds", subfield="bc", comparison_func=lambda x: x > 0, + target_kernel=target_kernel, ) if len(filtered_report_card) == 0: diff --git a/src/intelliperf/formulas/diagnose_only.py b/src/intelliperf/formulas/diagnose_only.py index b96b7533..43da42ce 100644 --- a/src/intelliperf/formulas/diagnose_only.py +++ b/src/intelliperf/formulas/diagnose_only.py @@ -64,7 +64,7 @@ def profile_pass(self): def instrument_pass(self): return super().instrument_pass() - def optimize_pass(self): + def optimize_pass(self, target_kernel: str = None): return super().optimize_pass() def compile_pass(self): diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 10da7c92..b836f4fa 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -235,7 +235,7 @@ def instrument_pass(self): self._application.build(instrumented=True) @abstractmethod - def optimize_pass(self): + def optimize_pass(self, target_kernel: str = None): """ Optimize the application based on the data collected from the instrumentation pass. """ @@ -329,6 +329,8 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole 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)}") + else: argument_name = kernel_args[i] logging.debug( @@ -513,7 +515,7 @@ def flatten_dict(d, parent_key="", sep="_"): return dict(items) -def filter_json_field(d, field, subfield=None, comparison_func=lambda x: True): +def filter_json_field(d, field, subfield=None, comparison_func=lambda x: True, target_kernel=None): """ Filters a list of dictionaries based on a comparison function applied to a specified field or subfield. @@ -527,9 +529,19 @@ def filter_json_field(d, field, subfield=None, comparison_func=lambda x: True): list: A list of dictionaries that satisfy the comparison function. """ if subfield is not None: - return [entry for entry in d if comparison_func(entry.get(field, {}).get(subfield, 0))] + return [ + entry + for entry in d + if comparison_func(entry.get(field, {}).get(subfield, 0)) + and (target_kernel is None or get_kernel_name(entry["kernel"]) == target_kernel) + ] else: - return [entry for entry in d if comparison_func(entry.get(field, 0))] + return [ + entry + for entry in d + if comparison_func(entry.get(field, 0)) + and (target_kernel is None or get_kernel_name(entry["kernel"]) == target_kernel) + ] def validate_arrays(arr1, arr2, tolerance): diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index 6baf222f..9e57628b 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -126,7 +126,7 @@ def instrument_pass(self) -> Result: error_report="The instrumentation is not implemented for memory access.", ) - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000) -> Result: + def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: """ Optimize the kernel to remove uncoalesced memory access via OpenAI API @@ -172,6 +172,7 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000) -> Res field=field, subfield=subfield, comparison_func=lambda x: x < peak_coal, + target_kernel=target_kernel, ) if len(filtered_report_card) == 0: diff --git a/src/intelliperf/formulas/swizzling.py b/src/intelliperf/formulas/swizzling.py new file mode 100644 index 00000000..68abc86c --- /dev/null +++ b/src/intelliperf/formulas/swizzling.py @@ -0,0 +1,532 @@ +################################################################################ +# MIT License + +# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. + +# 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: + +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. + +# 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. +################################################################################ + +import difflib +import json +import logging +import os +import stat +import sys + +import dspy + +from intelliperf.core.gpu_spec import GPUSpec +from intelliperf.core.llm import LLM +from intelliperf.formulas.formula_base import ( + Formula_Base, + Result, + filter_json_field, + get_kernel_name, +) +from intelliperf.utils.env import get_llm_api_key + + +class SwizzlingOptimization(dspy.Signature): + """Optimize GPU kernel code by applying a swizzling pattern to improve L2 cache locality.""" + + prompt = dspy.InputField( + desc="The user prompt containing the original code, memory analysis, and optimization history." + ) + + reason_why_old_was_slow = dspy.OutputField( + desc='JSON dictionary where keys are "iteration X" and values are sentences explaining why the swizzling pattern in that iteration was suboptimal.' + ) + summary_of_optimization = dspy.OutputField( + desc="An overview of the new code swizzling optimization that will be implemented." + ) + reason_why_new_should_be_better = dspy.OutputField( + desc="A comparison of the new optimization to the old optimizations, explaining why it should be better." + ) + result_code = dspy.OutputField( + desc="The full kernel code with the new swizzling optimization applied. This code should be complete and runnable." + ) + + +class swizzling(Formula_Base): + def __init__( + self, + name: str, + build_command: list, + instrument_command: list, + project_directory: str, + app_cmd: list, + top_n: int, + only_consider_top_kernel=False, + model: str = "gpt-4o", + provider: str = "openai", + in_place: bool = False, + unittest_command: str = None, + ): + super().__init__( + name, + build_command, + instrument_command, + project_directory, + app_cmd, + top_n, + model, + provider, + in_place, + unittest_command, + ) + + # This temp option allows us to toggle if we want a full or partial instrumentation report + self.only_consider_top_kernel = only_consider_top_kernel + self._instrumentation_results = None + self.current_kernel = None + self.current_args = None + self.current_kernel_signature = None + self.kernel_to_optimize = None + self.optimization_report = None + self.bottleneck_report = None + self.current_summary = None + self.previous_source_code = None + self.memory_analysis_output = None + self.success = False + + # New fields for logging + self.memory_analysis_prompt = None + self.optimization_prompt = None + + # New fields for iteration history tracking + self.iteration_history = [] # List of dicts with {iteration, diff, report, success} + self.current_iteration = 0 + self.memory_analysis_done = False + self.last_applied_diff = None + self.initial_source_code = None + + self.max_iterations = 10 + self.best_l2_improvement = 0.0 # Start at baseline (no improvement) + self.best_speedup = 1.0 # Start at 1.0x (no speedup) + self.best_diff = "" + self.best_iteration_report = "" + self.best_kernel_code = "" + self.best_optimization_results = None + self.l2_improvement_history = [] + self.gpu_spec = GPUSpec() + + # Removed local compute_diff; using Formula_Base.compute_diff instead + + def build_pass(self, validate_build_result=True) -> Result: + """ + Build the application and store the summary. + + Args: + validate_build_result (bool): Whether to validate the build result + + Returns: + Result: Build status and the output file path + """ + result = super().build(validate_build_result=validate_build_result) + if not result: + self.current_summary = result.error_report + return result + + def profile_pass(self) -> Result: + """ + Profile the application using guided-tuning and collect l2 hit rate data + + Returns: + Result: DataFrame containing the performance report card + """ + return super().profile_pass() + + def instrument_pass(self) -> Result: + """ + Instrument the application, targeting the kernels with the lowest l2 hit rate + + Returns: + Result: Instrumentation data containing the kernel name, arguments, lines, and file path as dict + """ + super().instrument_pass() + + return Result( + success=False, + asset=self._instrumentation_results, + error_report="The instrumentation is not implemented for swizzling.", + ) + + def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: + """ + Optimize the kernel to improve l2 hit rate through block swizzling via two-stage LLM approach + + Args: + temperature (float): Sampling temperature for OpenAI API + max_tokens (int): Maximum tokens for OpenAI API + + Returns: + Result: Optimized kernel as a file path + """ + super().optimize_pass() + llm_key = get_llm_api_key() + + # Increment iteration counter + self.current_iteration += 1 + + # First stage: Memory analysis system prompt (only run once) + analysis_system_prompt = ( + "You are a skilled GPU programmer with deep expertise in memory access patterns and cache locality. " + "You will analyze code to understand memory access patterns and locality opportunities, " + "but you will not modify any code. Focus on providing detailed, accurate insights about " + "memory access patterns that can be used to improve cache locality through block swizzling." + ) + + # Second stage: Optimization system prompt + optimization_system_prompt = ( + "You are a skilled GPU programmer specializing in block swizzling optimization. " + "Given a kernel and memory access analysis, you will implement swizzling to improve L2 cache locality. " + "Do not modify the kernel signature. Do not touch any other code, licenses, copyrights, or comments in the file. " + "If you remove the copyright, your solution will be rejected. " + "Do not include any markdown code blocks or text other than the code." + ) + + provider = self.provider + model = self.model + + # Only create analysis_llm if we haven't done memory analysis yet + if not self.memory_analysis_done: + analysis_llm = LLM( + api_key=llm_key, + system_prompt=analysis_system_prompt, + model=model, + provider=provider, + ) + + optimization_llm = LLM( + api_key=llm_key, + system_prompt=optimization_system_prompt, + model=model, + provider=provider, + ) + + kernel = None + kernel_file = None + + if self._instrumentation_results is None: + # Get the file from the results - look for kernels with low l2 hit rate + field = "l2" + subfield = "hr" + min_l2_hit_rate = 95 # Look for kernels with less than 95% l2 hit rate + filtered_report_card = filter_json_field( + self._initial_profiler_results, + field=field, + subfield=subfield, + comparison_func=lambda x: x < min_l2_hit_rate, + target_kernel=target_kernel, + ) + + if len(filtered_report_card) == 0: + return Result(success=False, error_report="No kernels with low l2 hit rate found.") + + logging.debug(f"Filtered Report Card:\n{json.dumps(filtered_report_card, indent=4)}") + + kernel = filtered_report_card[0]["kernel"] + files = filtered_report_card[0]["source"]["files"] + kernel_name = get_kernel_name(kernel) + + logging.debug(f"Kernel name: {kernel_name}") + kernel_file = None + for file in files: + if os.path.exists(file): + with open(file, "r") as f: + unoptimized_file_content = f.read() + if kernel_name in unoptimized_file_content: + kernel_file = file + break + if kernel_file is None: + logging.error(f"Kernel file not found for kernel {kernel}") + sys.exit(1) + else: + logging.debug(f"Kernel file found for kernel {kernel}: {kernel_file}") + + # Stage 1: Memory access pattern analysis (only run once) + if not self.memory_analysis_done: + with open(kernel_file, "r") as f: + initial_file_content = f.read() + self.initial_source_code = initial_file_content + analysis_prompt = ( + f"{self.initial_source_code}\n\n" + "I have this kernel and am trying to understand the memory access patterns of the kernel, and where there is memory locality that can be taken advantage of in the hardware cache. I will use this to swizzle the block id to better align the work so we have better cache locality.\n\n" + "I DO NOT want you to rewrite any code. I only want you to give me an overview for the memory access patterns and memory locality of the kernel. This will be used as context for future prompts that will take advantage of your insights. Make sure these insights on memory access patterns and locality between blocks in the kernel are accuracy and insightful so that I can actually take advantage of them to improve locality." + ) + self.memory_analysis_prompt = analysis_prompt + + self.bottleneck_report = ( + f"L2 Cache Locality Detection: IntelliPerf identified suboptimal L2 cache hit rate " + f"in kernel `{kernel_name}`. Poor cache locality occurs when " + f"blocks accessing related memory are scheduled to different XCDs with separate L2 caches, " + f"reducing overall cache effectiveness." + ) + + # Stage 1: Get memory access analysis + try: + logging.debug(f"Analysis prompt: {analysis_prompt}") + self.memory_analysis_output = analysis_llm.ask( + analysis_prompt, + signature="prompt: str -> memory_analysis_output: str", + answer_type="memory_analysis_output", + ) + self.memory_analysis_output = self.memory_analysis_output.strip() + + logging.debug(f"Memory analysis output: {self.memory_analysis_output}") + + self.memory_analysis_done = True + except Exception as e: + logging.error(f"Failed to get memory analysis - {str(e)}") + return Result(success=False, error_report=f"Failed to get memory analysis - {str(e)}") + + history_prompt_part = "" + if self.iteration_history: + history_prompt_part += "Here is the history of previous optimization attempts (Note that YOU ARE NOT ALLOWED TO REIMPLEMENT THE SAME SWIZZLING PATTERN):\n\n" + for item in self.iteration_history: + history_prompt_part += f"--- Iteration {item['iteration']} ---\n" + history_prompt_part += f"Applied diff:\n{item['diff']}\n" + history_prompt_part += f"Profiling report:\n{item['report']}\n\n" + + # Stage 2: Swizzling optimization + optimization_prompt = ( + f"The original code is:\n\n {self.initial_source_code}\n\n" + f"The memory analysis is:\n\n {self.memory_analysis_output}\n\n" + f"{history_prompt_part}" + "Pay special attention to the swizzling pattern in the diff. If you see a swizzling pattern in the diff, do not reimplement it. Instead, try to implement an completely new approach to swizzling." + "On the MI300x GPU there are multiple XCDs, and each XCD has its own L2 cache. So that blocks on the same XCD that access the same memory will likely hit in the shared L2 cache and thus improve the L2 hit rate of the program. For this reason, blocks that share the same data should be scheduled to the same XCD. Your task is to find the swizzling formulation such that blocks that access the same memory will be scheduled to the same XCD.\n\n" + "MI300X architecture specification\n\n" + f"The GPU contains {self.gpu_spec.get_num_xcds()} XCDs.\n\n" + f"Each XCD has its own L2 cache of {self.gpu_spec.get_l2_cache_size_per_xcd()} KB.\n\n" + f"Each XCD has {self.gpu_spec.get_num_cus_per_xcd()} CUs, and blocks are assigned to CUs.\n\n" + "We want to maximize utilization by assigning an equal number of blocks to each XCD.\n\n" + f"{self.gpu_spec.get_hip_runtime_scheduling_info()}" + "Swizzling goal\n\n" + "Recompute the block index with the old block index, number of XCDs on the GPU, and total number of blocks in the program so that:\n\n" + "Blocks that share the same data map to the same XCD until that XCD's share is filled.\n\n" + "Work remains evenly balanced across XCDs.\n\n" + "We want to understand how the blocks are strided by the round robin scheduler. Some question that you might want to ask (but might not necessarily be relevant) are: How do you calculate the XCD id that the block was originally mapped to? How many blocks are in each XCD/ How do you understand the stride of block indexes. If we have to round robin for multiple iterations, how do we calculate the number of iterations that the block index was assigned on? How can we use this to make an offset for reassigning the block index.\n\n" + "There are potentially many more question that you might want to ask when understanding how to best swizzle the kernel to take advantage of locality. To be very clear, the optimal swizzling pattern will change by algorithm. Different algorithms reuse data differently, and thus the blocks that should share the same L2 cache will change by different algorithms based on memory access patterns. I want you to deeply understand how to do this for the specific algorithm we are working on.\n\n" + "I want you to consider the swizzling pattern step by step and then put everything together in the formula.\n\n" + "Task\n\n" + f"number of XCDs = {self.gpu_spec.get_num_xcds()} in this hardware architecture. In the case of this program, number of blocks is equal to number of SMS, so you can directly use that argument. Make sure you do not change the parameters in the kernel function, as this will break the code. The function signature must stay exactly the same or the code will fail.\n\n" + "Propose a swizzling pattern as one or a few lines of code inside the kernel that reassigns the block index. For HIP kernels, you must still eventually assign threadId. For the HIP kernels, also make sure that you use the new swizzled block ids for all thread id computation. For Triton kernels, you must still eventually assign pid. Rewrite the code of the entire kernel without putting in any placeholders. I want to be able to take the code, copy it into a new file, and run it on the testbench without any extra work. Again, make sure to not change the kernel function signature and only add new swizzling lines within the kernel using the available parameters.\n\n" + "EXTREMELY IMPORTANT - Do not include any markdown code blocks or text other than the code. DO NOT start the code with 'python'. I want you to straight directly output the code. I want to be able to copy and paste the code into a new file and run it on the testbench without any extra work. DO NOT REMOVE ANY CODE. DO NOT MODIFY ANY HOST SIDE CODE.\n\n" + "EXTREMELY IMPORTANT - Make sure to not change the kernel function signature. Do not add any new parameters to the kernel function. Do not change the return type of the kernel function. Do not change the name of the kernel function. Do not change the arguments of the kernel function. Do not change the return type of the kernel function. Do not change the name of the kernel function. Do not change the arguments of the kernel function. Do not change the return type of the kernel function. Do not change the name of the kernel function. Do not change the arguments of the kernel function.\n\n" + "EXTREMELY IMPORTANT - I always want the original pid to be written to a variable called pid and ending in a variable called pid. If we have a 2D grid of pids, they must be called pid_x and pid_y, and so on. It is very important that you name the variables by this format and write the whole code based around these variable names so that it runs successfully.\n\n" + "EXTREMELY IMPORTANT - Make sure your output is in the correct format. The fields are reason_why_old_was_slow, summary_of_optimization, reason_why_new_should_be_better, result_code.\n\n" + "**OPTIMIZATION GOAL** - You have not reached the maximum performance yet. DO NOT REIMPLEMENT A PREVIOUS SWIZZLING PATTERN. If you have previously tried an approach and it is shown in the diff, do not reimplement it.\n\n" + ) + + if self.current_summary is not None: + optimization_prompt += f"\n\nThe current summary is: {self.current_summary}" + cur_diff = self.compute_diff([kernel_file]) + optimization_prompt += f"\nThe diff between the current and initial code is: {cur_diff}" + else: + pass + + if kernel is None: + return Result(success=False, error_report="Failed to extract the kernel name.") + if kernel_file is None: + return Result(success=False, error_report="Failed to extract the kernel file path.") + + logging.debug(f"Optimization prompt: {optimization_prompt}") + + self.current_kernel = kernel.split("(")[0] + # self.current_args = kernel.split("(")[1].split(")")[0].split(",") + self.current_kernel_signature = kernel + + self.current_kernel_files = [kernel_file] + try: + with open(kernel_file, "r") as f: + code_before_opt = f.read() + + response = optimization_llm.ask(optimization_prompt, signature=SwizzlingOptimization) + optimized_file_content = response.result_code.strip() + + # Strip markdown code blocks if present using Formula_Base helper + optimized_file_content = self.postprocess_llm_code(optimized_file_content) + + # If this is a Python kernel file, ensure it starts with a shebang + if kernel_file.endswith(".py") and not optimized_file_content.startswith("#!"): + optimized_file_content = "#!/usr/bin/env python\n" + optimized_file_content + + diff = difflib.unified_diff( + code_before_opt.splitlines(True), + optimized_file_content.splitlines(True), + fromfile=f"a/{os.path.basename(kernel_file)}", + tofile=f"b/{os.path.basename(kernel_file)}", + ) + self.last_applied_diff = "".join(list(diff)) + with open(kernel_file, "w") as f: + f.write(optimized_file_content) + + # Mark the optimized file as executable (chmod +x) + try: + mode = os.stat(kernel_file).st_mode + os.chmod(kernel_file, mode | stat.S_IXUSR | stat.S_IXGRP | stat.S_IXOTH) + except Exception as e: + logging.debug(f"Failed to chmod +x on {kernel_file}: {e}") + + logging.debug(f"Optimized file content: {optimized_file_content}") + return Result( + success=True, + asset={ + "optimized_code_path": kernel_file, + "optimized_code_string": optimized_file_content, + }, + ) + except Exception as e: + logging.error(f"An unexpected error occurred - {str(e)}") + return Result(success=False, error_report=f"An unexpected error occurred - {str(e)}") + + def compiler_pass(self) -> Result: + """ + Compile the application + + Returns: + Result: Compilation status and the output file path + """ + return super().compile_pass() + + def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) -> Result: + """ + Validate the optimized kernel by running the provided unittest_command. + + If no unittest_command is provided, skip validation (treat as success). + """ + if not self.unittest_command: + return Result(success=True, asset={"log": "No unittest_command provided; skipping correctness validation."}) + + success, output = self._application.run_unit_test() + if not success: + return Result(success=False, error_report=f"Unit test validation failed. Output:\n{output}") + return Result(success=True, asset={"log": output}) + + def performance_validation_pass(self) -> Result: + unoptimized_results = filter_json_field( + self._initial_profiler_results, + field="kernel", + comparison_func=lambda x: x == self.current_kernel_signature, + ) + + unoptimized_time = unoptimized_results[0]["durations"]["ns"] + unoptimized_l2_hit_rate = unoptimized_results[0]["l2"]["hr"] + kernel = unoptimized_results[0]["kernel"] + + # Profile the optimized application + self._optimization_results = self._application.profile(top_n=self.top_n) + + optimized_results = filter_json_field( + self._optimization_results, + field="kernel", + comparison_func=lambda x: x == kernel, + ) + + optimized_time = optimized_results[0]["durations"]["ns"] + optimized_l2_hit_rate = optimized_results[0]["l2"]["hr"] + + success = optimized_l2_hit_rate > unoptimized_l2_hit_rate + speedup = unoptimized_time / optimized_time + l2_improvement = optimized_l2_hit_rate - unoptimized_l2_hit_rate + + self.optimization_report = "" + self.l2_improvement_history.append(l2_improvement) + + # Format the L2 cache improvement message + if l2_improvement > 0: + self.optimization_report += ( + f"L2 Cache Locality Improvement: Successfully improved L2 cache hit rate by " + f"{l2_improvement:.2f} percentage points. " + f"Hit rate increased from {unoptimized_l2_hit_rate:.1f}% to {optimized_l2_hit_rate:.1f}% " + f"(higher percentages indicate better cache locality through improved block swizzling). " + ) + else: + self.optimization_report += ( + f"L2 Cache Locality Degradation: L2 cache hit rate decreased by " + f"{abs(l2_improvement):.2f} percentage points. " + f"Hit rate decreased from {unoptimized_l2_hit_rate:.1f}% to {optimized_l2_hit_rate:.1f}% " + f"(lower percentages indicate worse cache locality). " + ) + + # 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"({(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"({(1 / speedup - 1) * 100:.1f}% slower)." + ) + + self.iteration_history.append( + { + "iteration": self.current_iteration, + "diff": self.last_applied_diff, + "report": self.optimization_report, + "success": success, + } + ) + + # Update best if speedup improved AND L2 is better than reference (baseline) + is_better = speedup > self.best_speedup and l2_improvement > 0 + + if is_better: + self.best_l2_improvement = l2_improvement + self.best_speedup = speedup + self.best_diff = self.last_applied_diff + self.best_iteration_report = self.optimization_report + self.best_optimization_results = self._optimization_results + with open(self.current_kernel_files[0], "r") as f: + self.best_kernel_code = f.read() + # Mark as successful if we achieved any improvement + if l2_improvement > 0 or speedup > 1.0: + self.success = True + if self.current_iteration < self.max_iterations: + self.current_summary = self.optimization_report + # Always return success=False to continue iterating + return Result(success=False, error_report=self.best_iteration_report) + + return Result(success=True, asset={"log": self.best_iteration_report}) + + def write_results(self, output_file: str = None): + """ + Writes the results to the output file. + """ + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report + + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) + + super().write_results( + output_file=output_file, + additional_results={"formula": "swizzling", "success": self.success}, + ) + + def summarize_previous_passes(self): + """ + Summarizes the results of the previous passes for future prompts. + """ + pass From 934f9ca8889e3c936368f4aafe39b8500330a69b Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Wed, 15 Oct 2025 02:32:33 -0700 Subject: [PATCH 05/14] Add histroy to all formulas (#148) Co-authored-by: github-actions[bot] --- .../workflows/scripts/check_test_results.sh | 10 + .github/workflows/scripts/ci_tests.sh | 157 ++++- examples/CMakeLists.txt | 3 +- src/intelliperf/__main__.py | 7 +- src/intelliperf/core/llm.py | 50 +- src/intelliperf/core/logger.py | 88 ++- src/intelliperf/formulas/atomic_contention.py | 572 ++++++++++++++---- src/intelliperf/formulas/bank_conflict.py | 492 +++++++++++++-- src/intelliperf/formulas/formula_base.py | 222 ++++++- src/intelliperf/formulas/memory_access.py | 491 +++++++++++++-- src/intelliperf/formulas/swizzling.py | 60 +- 11 files changed, 1879 insertions(+), 273 deletions(-) diff --git a/.github/workflows/scripts/check_test_results.sh b/.github/workflows/scripts/check_test_results.sh index 00821a0a..e9d257e4 100755 --- a/.github/workflows/scripts/check_test_results.sh +++ b/.github/workflows/scripts/check_test_results.sh @@ -1,6 +1,14 @@ #!/bin/bash set -e +# Check if jq is installed +if ! command -v jq &> /dev/null; then + echo "❌ Error: jq is not installed" + echo " Please install jq to parse JSON test results" + echo " Install with: apt-get install jq (Ubuntu/Debian) or yum install jq (RHEL/CentOS)" + exit 1 +fi + # Check test results for success echo 'Checking test results...' results_dir=~/intelliperf_results @@ -9,6 +17,8 @@ results_dir=~/intelliperf_results check_test_result() { local file="$1" local test_name="$2" + echo "Checking test result for $test_name" + echo "File: $file" if [ -f "$file" ]; then if jq -e '.success == true' "$file" >/dev/null 2>&1; then echo "βœ… $test_name: PASSED" diff --git a/.github/workflows/scripts/ci_tests.sh b/.github/workflows/scripts/ci_tests.sh index 34848ea7..d4d03661 100755 --- a/.github/workflows/scripts/ci_tests.sh +++ b/.github/workflows/scripts/ci_tests.sh @@ -1,35 +1,142 @@ #!/bin/bash set -e -# Run examples and store outputs +#=============================================================================== +# IntelliPerf CI Test Suite +# This script runs various IntelliPerf examples and logs their output +#=============================================================================== + +echo '================================' echo 'Running IntelliPerf examples...' +echo '================================' +# Setup results directory results_dir=~/intelliperf_results -rm -rf $results_dir -mkdir -p $results_dir - +rm -rf "$results_dir" +mkdir -p "$results_dir" +# Configuration provider="openrouter" model="openai/gpt-4o" -# Formulas -intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --formula=memoryAccess -o $results_dir/memory_access_output.json -- ./build/access_pattern/uncoalesced || true -intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --formula=bankConflict -o $results_dir/bank_conflict_output.json -- ./build/bank_conflict/matrix_transpose 1024 1024 || true -intelliperf -vvv --project_directory=./examples --provider $provider --model $model --build_command="./scripts/build_examples.sh -c" --instrument_command="./scripts/build_examples.sh -i -c" --formula=atomicContention -o $results_dir/atomic_contention_output.json -- ./build/contention/reduction || true -intelliperf -vvv --project_directory=./examples --provider $provider --model $model --formula=swizzling --project_directory="./examples" --unittest_command="triton/gemm_runner.py --validate" -o $results_dir/swizzling_output.json -- ./triton/gemm_runner.py || true -# Diagnose Only -intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_hip_uncoalesced.json -- ./examples/build/access_pattern/uncoalesced -intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_torch_add.json -- ./examples/torch/add.py -TRITON_DISABLE_LINE_INFO=0 intelliperf -vvv --formula=diagnoseOnly -o $results_dir/diagnose_only_triton_reduce.json -- ./examples/triton/reduce.py - -# Display output files -echo 'Memory Access Output:' -cat $results_dir/memory_access_output.json || echo "File not found" -echo 'Bank Conflict Output:' -cat $results_dir/bank_conflict_output.json || echo "File not found" -echo 'Atomic Contention Output:' -cat $results_dir/atomic_contention_output.json || echo "File not found" -echo 'Diagnose Only Output:' -cat $results_dir/diagnose_only_hip_uncoalesced.json || echo "File not found" -cat $results_dir/diagnose_only_torch_add.json || echo "File not found" -cat $results_dir/diagnose_only_triton_reduce.json || echo "File not found" +#=============================================================================== +# Formula-based Tests +#=============================================================================== + +echo "" +echo "[1/7] Running Memory Access formula test..." +intelliperf -vvv \ + --project_directory=./examples \ + --provider "$provider" \ + --model "$model" \ + --build_command="./scripts/build_examples.sh -c" \ + --formula=memoryAccess \ + -o "$results_dir/memory_access_output.json" \ + --trace_path "$results_dir/memory_access" \ + -- ./build/access_pattern/uncoalesced 2>&1 | tee "$results_dir/memory_access_output.log" || true + +echo "" +echo "[2/7] Running Bank Conflict formula test..." +intelliperf -vvv \ + --project_directory=./examples \ + --provider "$provider" \ + --model "$model" \ + --build_command="./scripts/build_examples.sh -c" \ + --formula=bankConflict \ + -o "$results_dir/bank_conflict_output.json" \ + --trace_path "$results_dir/bank_conflict" \ + -- ./build/bank_conflict/matrix_transpose 1024 1024 2>&1 | tee "$results_dir/bank_conflict_output.log" || true + +echo "" +echo "[3/7] Running Atomic Contention formula test..." +intelliperf -vvv \ + --project_directory=./examples \ + --provider "$provider" \ + --model "$model" \ + --build_command="./scripts/build_examples.sh -c" \ + --instrument_command="./scripts/build_examples.sh -i -c" \ + --formula=atomicContention \ + -o "$results_dir/atomic_contention_output.json" \ + --trace_path "$results_dir/atomic_contention" \ + -- ./build/contention/reduction 2>&1 | tee "$results_dir/atomic_contention_output.log" || true + +echo "" +echo "[4/7] Running Swizzling formula test..." +intelliperf -vvv \ + --project_directory=./examples \ + --provider "$provider" \ + --model "$model" \ + --formula=swizzling \ + --unittest_command="triton/gemm_runner.py --validate" \ + -o "$results_dir/swizzling_output.json" \ + -- ./triton/gemm_runner.py 2>&1 | tee "$results_dir/swizzling_output.log" || true + +#=============================================================================== +# Diagnose-Only Tests +#=============================================================================== + +echo "" +echo "[5/7] Running Diagnose Only test (HIP uncoalesced)..." +intelliperf -vvv \ + --formula=diagnoseOnly \ + -o "$results_dir/diagnose_only_hip_uncoalesced.json" \ + --trace_path "$results_dir/diagnose_only_hip_uncoalesced" \ + -- ./examples/build/access_pattern/uncoalesced 2>&1 | tee "$results_dir/diagnose_only_hip_uncoalesced.log" + +echo "" +echo "[6/7] Running Diagnose Only test (Torch add)..." +intelliperf -vvv \ + --formula=diagnoseOnly \ + -o "$results_dir/diagnose_only_torch_add.json" \ + --trace_path "$results_dir/diagnose_only_torch_add" \ + -- ./examples/torch/add.py 2>&1 | tee "$results_dir/diagnose_only_torch_add.log" + +echo "" +echo "[7/7] Running Diagnose Only test (Triton reduce)..." +TRITON_DISABLE_LINE_INFO=0 intelliperf -vvv \ + --formula=diagnoseOnly \ + -o "$results_dir/diagnose_only_triton_reduce.json" \ + --trace_path "$results_dir/diagnose_only_triton_reduce" \ + -- ./examples/triton/reduce.py 2>&1 | tee "$results_dir/diagnose_only_triton_reduce.log" + +#=============================================================================== +# Display Results Summary +#=============================================================================== + +echo "" +echo '=======================================' +echo 'Test Results Summary' +echo '=======================================' + +echo "" +echo '[Memory Access Output]' +cat "$results_dir/memory_access_output.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '[Bank Conflict Output]' +cat "$results_dir/bank_conflict_output.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '[Atomic Contention Output]' +cat "$results_dir/atomic_contention_output.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '[Swizzling Output]' +cat "$results_dir/swizzling_output.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '[Diagnose Only - HIP Uncoalesced]' +cat "$results_dir/diagnose_only_hip_uncoalesced.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '[Diagnose Only - Torch Add]' +cat "$results_dir/diagnose_only_torch_add.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '[Diagnose Only - Triton Reduce]' +cat "$results_dir/diagnose_only_triton_reduce.json" 2>/dev/null || echo " ⚠ File not found" + +echo "" +echo '=======================================' +echo 'All logs saved to:' "$results_dir/*.log" +echo '=======================================' diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 16a0d784..b3c34db1 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -44,7 +44,8 @@ function(add_example name source_file) endif() # Must include debug information to get line numbers endif() - target_compile_options(${name} PRIVATE -g) + target_compile_options(${name} PRIVATE -g) + target_compile_options(${name} PRIVATE -gline-tables-only) endfunction() add_subdirectory(bank_conflict) diff --git a/src/intelliperf/__main__.py b/src/intelliperf/__main__.py index bb731769..bf2d0926 100644 --- a/src/intelliperf/__main__.py +++ b/src/intelliperf/__main__.py @@ -240,6 +240,8 @@ def main(): sys.exit(1) + num_attempts = 0 if args.formula == "diagnoseOnly" else args.num_attempts + optimizer_args = { "name": generated_name, "build_command": args.build_command, @@ -251,6 +253,7 @@ def main(): "provider": args.provider, "in_place": args.in_place, "unittest_command": args.unittest_command, + "num_attempts": num_attempts, } optimizer = formula(**optimizer_args) @@ -324,7 +327,9 @@ def flush_logs_if_enabled(): import sys try: - if args.formula == "diagnoseOnly" or performance_result: + # Always write results if we ran any iterations (even if they all "failed" to continue iterating) + # For diagnoseOnly or if we have any performance result, write output + if args.formula == "diagnoseOnly" or performance_result is not None: # Flush logger if tracing is enabled if hasattr(optimizer, "get_logger") and args.trace_path: logger = optimizer.get_logger() diff --git a/src/intelliperf/core/llm.py b/src/intelliperf/core/llm.py index 0eece144..b4df30bd 100644 --- a/src/intelliperf/core/llm.py +++ b/src/intelliperf/core/llm.py @@ -99,13 +99,34 @@ def __init__( def ask( self, - user_prompt: str, + user_prompt: str = "", signature="prompt: str -> optimized_code: str", answer_type: str = "optimized_code", record_meta: str = None, + **input_kwargs, ): + """ + Ask the LLM a question using DSPy signatures. + + Args: + user_prompt: For simple string-based signatures, the prompt text + signature: DSPy signature (string or Signature class) + answer_type: Which field to extract from response (for string signatures) + record_meta: Metadata for logging + **input_kwargs: For complex signatures with multiple inputs (e.g., kernel_code, history, etc.) + """ # Log the LLM interaction start if self.logger: + # For input_kwargs, log the actual values + logged_inputs = {} + if input_kwargs: + for key, value in input_kwargs.items(): + # For history object, extract the messages list + if key == "history" and hasattr(value, "messages"): + logged_inputs[key] = value.messages + else: + logged_inputs[key] = value + self.logger.record( "llm_call_start", { @@ -116,6 +137,7 @@ def ask( "record_meta": record_meta, "signature": str(signature), "answer_type": answer_type, + "input_kwargs": logged_inputs if input_kwargs else None, }, ) @@ -152,7 +174,14 @@ def ask( else: # DSPy path dspy.context(description=self.system_prompt) chain = dspy.ChainOfThought(signature) - ct_response = chain(prompt=user_prompt) + + # Determine how to call the chain + if input_kwargs: + # Complex signature with multiple inputs + ct_response = chain(**input_kwargs) + else: + # Simple signature with just prompt + ct_response = chain(prompt=user_prompt) # Try to capture reasoning if available (not returned) reasoning = getattr(ct_response, "reasoning", None) @@ -171,16 +200,8 @@ def ask( "record_meta": record_meta, "signature": str(signature), "answer_type": answer_type, + "response": response_content, } - try: - if isinstance(response_content, str): - log_payload["response"] = response_content - log_payload["response_length"] = len(response_content) - else: - # Best-effort: record fields present on prediction object - log_payload["response_fields"] = list(getattr(ct_response, "__dict__", {}).keys()) - except Exception: - pass if reasoning: log_payload["reasoning"] = reasoning log_payload["reasoning_type"] = "chain_of_thought" @@ -193,7 +214,12 @@ def ask( error_type = type(e).__name__ if self.logger: self.logger.record( - "llm_call_error", {"error": error_message, "error_type": error_type, "record_meta": record_meta} + "llm_call_error", + { + "error": error_message, + "error_type": error_type, + "record_meta": record_meta, + }, ) print(f"ERROR: {error_message}") sys.exit(1) diff --git a/src/intelliperf/core/logger.py b/src/intelliperf/core/logger.py index a5ebc215..d52f9170 100644 --- a/src/intelliperf/core/logger.py +++ b/src/intelliperf/core/logger.py @@ -45,7 +45,14 @@ def __init__(self, run_name: str = None): self.start_time = time.time() # Record run start - self.record("run_start", {"run_id": self.run_id, "run_name": self.run_name, "timestamp": self.start_time}) + self.record( + "run_start", + { + "run_id": self.run_id, + "run_name": self.run_name, + "timestamp": self.start_time, + }, + ) def __del__(self): """Destructor to ensure logs are flushed when Logger object is destroyed""" @@ -64,12 +71,89 @@ def record(self, event_type: str, data: Dict[str, Any]) -> None: event_type: Type of event (e.g., "llm_call", "optimization_pass", "validation_result") data: Event data dictionary """ - entry = {"timestamp": time.time(), "type": event_type, "data": data} + # Serialize complex objects in data + serialized_data = self._serialize_data(data) + entry = {"timestamp": time.time(), "type": event_type, "data": serialized_data} self.buffer.append(entry) # Also log to console for immediate visibility logging.debug(f"Logger: {event_type} - {data}") + def _serialize_data(self, data: Dict[str, Any]) -> Dict[str, Any]: + """ + Recursively serialize data, converting complex objects to dictionaries. + + Args: + data: Data dictionary to serialize + + Returns: + Serialized data dictionary + """ + serialized = {} + for key, value in data.items(): + serialized[key] = self._serialize_value(value) + + # Automatically add diff_lines when diff is present + if "diff" in serialized and "diff_lines" not in serialized: + diff_value = serialized["diff"] + if isinstance(diff_value, str): + serialized["diff_lines"] = diff_value.split("\n") if diff_value else [] + + return serialized + + def _serialize_value(self, value: Any) -> Any: + """ + Serialize a single value, handling various types. + + Args: + value: Value to serialize + + Returns: + Serialized value + """ + # Handle basic types + if value is None or isinstance(value, (str, int, float, bool)): + return value + + # Handle lists and tuples + if isinstance(value, (list, tuple)): + return [self._serialize_value(item) for item in value] + + # Handle dictionaries + if isinstance(value, dict): + result = {k: self._serialize_value(v) for k, v in value.items()} + # Add diff_lines if diff exists and diff_lines doesn't + if "diff" in result and "diff_lines" not in result: + diff_value = result["diff"] + if isinstance(diff_value, str): + result["diff_lines"] = diff_value.split("\n") if diff_value else [] + return result + + # Handle objects with __dict__ + if hasattr(value, "__dict__"): + obj_dict = {} + for attr, attr_value in value.__dict__.items(): + if not attr.startswith("_"): + obj_dict[attr] = self._serialize_value(attr_value) + return obj_dict + + # Handle objects with accessible attributes via dir() + if hasattr(value, "__class__"): + obj_dict = {} + for attr in dir(value): + if not attr.startswith("_"): + try: + attr_value = getattr(value, attr) + if not callable(attr_value): + obj_dict[attr] = self._serialize_value(attr_value) + except Exception: + pass + if obj_dict: + return obj_dict + + # Fall back to string representation + return str(value) + def get_buffer(self) -> List[Dict[str, Any]]: """ Get current buffer - for debugging/fallback. diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index 8bc8f68a..0ec0b732 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -26,14 +26,197 @@ import logging import os -from intelliperf.core.llm import LLM +import dspy + from intelliperf.formulas.formula_base import ( Formula_Base, + OptimizationTracker, Result, filter_json_field, get_kernel_name, ) -from intelliperf.utils.env import get_llm_api_key + + +class AtomicContentionOptimization(dspy.Signature): + """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.") + + code_preservation_rules = dspy.InputField( + desc="""CRITICAL: These rules MUST be followed to avoid compilation errors: + +1. **NEVER modify, abbreviate, or remove copyright/license headers** + - Copy the ENTIRE copyright header exactly as written, character by character + - DO NOT use '...' or any ellipsis inside comments - this creates unterminated comment errors + - If you see: /* Copyright (c) 2025... */ you MUST write the full text + +2. **NEVER use '...' or ellipsis in C/C++ comments** + - The sequence '...' does NOT close a /* */ comment block + - This causes: error: unterminated /* comment + - Either copy the full comment text or remove it entirely (but keep copyright!) + +3. **Preserve the EXACT kernel signature** + - Do not change: function name, parameter types, parameter names, or return type + - The kernel must be callable with the same arguments as before""" + ) + + output_format_rules = dspy.InputField( + desc="""CRITICAL: Output format requirements: + +1. **No markdown code blocks or formatting** + - Do NOT include: ```cpp, ```c, ```hip, or ``` markers + - Do NOT include any explanatory text before or after the code + +2. **Output ONLY the complete, compilable source code** + - Include all necessary #include statements + - Include the full copyright header (copied verbatim) + - Include all function implementations + - The output must be ready to write directly to a .hip file + +3. **Preserve all existing comments and licenses** + - Copy them exactly as they appear in the original code""" + ) + + problem_description = dspy.InputField( + desc="Description of the atomic contention issue detected, including specific atomic operations with high latency and performance impact." + ) + + # Baseline metrics for reference + baseline_atomic_latency = dspy.InputField( + desc="Baseline average atomic latency in cycles. Indicates contention level - higher latency means more threads competing for atomic operations. This is the initial unoptimized value before any optimization attempts." + ) + baseline_time_ms = dspy.InputField( + desc="Baseline kernel execution time in milliseconds. This is the initial unoptimized kernel runtime before any optimization attempts." + ) + + # History of all previous attempts + history = dspy.History = dspy.InputField( + desc="Complete history of previous optimization attempts, including the code changes (diffs), before/after atomic latency values, speedup ratios, and whether each attempt improved or regressed performance. Use this to avoid repeating failed approaches and build on successful patterns." + ) + + previous_failures = dspy.InputField( + desc="CRITICAL: Analysis of previous failed attempts. DO NOT repeat these mistakes: compilation errors (unterminated comments, missing semicolons), correctness failures (changed semantics, wrong outputs), or performance regressions (same approach with minor variations). Learn from these failures and try fundamentally different optimization strategies." + ) + + latest_diff = dspy.InputField( + desc="The most recent code change that was attempted. If this resulted in a failure, you MUST try a completely different approach. Do not make minor variations of the same optimization - innovate with a new strategy." + ) + + # Low-level optimization techniques + optimization_techniques = dspy.InputField( + desc="""Available low-level atomic contention optimization techniques to consider: + +1. **Non-temporal loads**: Treat read-once data as streaming, avoid cache pollution. + Example: float x = __builtin_nontemporal_load(&input[i]); + +2. **Warp-wide cooperative loads** using shuffles: + float val = __shfl_sync(0xffffffff, local_data, 0); + +3. **Prefetching into registers**: + float prefetch = input[i + 1]; // Use later + +4. **Manual register tiling with unrolled loops**: + float tile[4]; + #pragma unroll + for (int i = 0; i < 4; ++i) tile[i] = input[idx + i]; + +5. **Double-buffering** for overlapping memory + compute: + load_tile(buf_a); + load_tile(buf_b); compute(buf_a); swap(); + +6. **Structure-of-Arrays transformation** to align thread memory access: + float val = soa.x[threadIdx.x]; // Coalesced + +7. **Aligned vectorized memory access** using float4: + float4 x = reinterpret_cast(input)[idx];""" + ) + + amd_specific_optimizations = dspy.InputField( + desc="""AMD-specific optimizations for CDNA GPUs (MI250X, MI300X): + +**AMD MFMA (Matrix Fused Multiply-Add) intrinsics**: +These instructions execute on a **wavefront-wide** basis (64 threads), using per-lane vector registers to load parts of matrices A, B, C. + +**MFMA Intrinsic Syntax:** + d = __builtin_amdgcn_mfma__(a, b, c, cbsz, abid, blgp); + +**Parameters:** +- CDfmt: format of C and D (e.g., f32, fp32) +- ABfmt: format of A and B (e.g., f16, bf16, i8) +- M, N, K: matrix tile dimensions +- a, b, c: registers or scalars from matrices A, B, C +- d: output register of resulting matrix tile D +- cbsz: broadcast size (0=no broadcast, 1=2-wide broadcast, etc.) +- abid: A-matrix block to broadcast from (used with cbsz) +- blgp: B-matrix swizzle pattern (0=normal, 1=lanes 0-31β†’32-63, 2=lanes 32-63β†’0-31, 3=rotate down by 16, 4-7=group-wide broadcast) + +**Example (FP32, 16x16x4 GEMM):** +```cpp +__global__ void sgemm_16x16x4(const float *A, const float *B, float *D) { + using float4 = __attribute__((__vector_size__(4 * sizeof(float)))) float; + float4 d = {0}; + + int mk = threadIdx.y + 4 * threadIdx.x; + int kn = threadIdx.x + 16 * threadIdx.y; + + float amk = A[mk]; + float bkn = B[kn]; + d = __builtin_amdgcn_mfma_f32_16x16x4f32(amk, bkn, d, 0, 0, 0); + + for (int i = 0; i < 4; ++i) + D[threadIdx.x + i * 16 + threadIdx.y * 4 * 16] = d[i]; +} +``` +Launch with: dim3 block(16, 4); dim3 grid(1, 1); + +**Performance (CDNA2/CDNA3 - MI250X/MI300X):** +- FP32 16x16x4: 256 flops/cycle/CU +- FP16/BF16 16x16x16: 1024 flops/cycle/CU +- INT8 16x16x16: 1024 flops/cycle/CU +- FP64 4x4x4: 128-256 flops/cycle/CU + +**Key Benefits:** +- MFMA offers 2-4Γ— throughput over vector FMAs +- Use wavefront-aligned loads and register blocking to match instruction layout +- Avoid using warp-level primitives with the _sync suffix (e.g., __shfl_sync, __ballot_sync) and use the non-sync variants instead for broader compatibility +- We are optimizing for CDNA GPUs (MI300X) so target this architecture""" + ) + + optimization_categories = dspy.InputField( + desc="""Categories of optimization strategies to explore: + +1. **Atomic Reduction Strategies:** + - Replace atomics with thread-local accumulation + final reduction + - Use hierarchical reduction (warp-level β†’ block-level β†’ global) + - Batch atomic operations to reduce frequency + - Use lock-free data structures when possible + +2. **Memory Access Pattern Changes:** + - Non-temporal loads for streaming data (__ldg, __builtin_nontemporal_load) + - Warp-wide cooperative loads with shuffles (__shfl, __shfl_down) + - Manual register tiling and double-buffering + - Structure-of-Arrays transformations + +3. **Advanced Vectorization:** + - Vectorized memory access patterns (float4, int4) + - SIMD-friendly loop unrolling + - Memory coalescing with vector types + +4. **AMD-Specific Optimizations:** + - AMD MFMA (Matrix Fused Multiply-Add) intrinsics for CDNA GPUs (MI250X, MI300X) + - Wavefront-level optimizations (64-thread blocks) + - AMD-specific memory hierarchy optimizations + +5. **Algorithmic Changes:** + - Different loop ordering strategies + - Alternative memory layout schemes + - New data structure organizations + - Wavefront/warp specialization and/or workgroup/block specialization""" + ) + + optimized_code = dspy.OutputField( + desc="The complete, runnable optimized kernel code with reduced atomic contention. Must preserve the exact kernel signature, all comments, licenses, and copyright notices. The code should reduce atomic operations or reorganize computation to minimize contention." + ) class atomic_contention(Formula_Base): @@ -73,8 +256,28 @@ def __init__( self.current_summary = None self.previous_source_code = None self.success = False - self.optimization_attempts = [] # Track optimization strategies attempted - self.iteration_count = 0 # Track current iteration + + # Initialize optimization tracker + # 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", + maximize=True, + before_metric="unoptimized_lat", + after_metric="optimized_lat", + ) + + # Store baseline metrics (set during first profiling) + 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: """ @@ -85,13 +288,14 @@ def profile_pass(self) -> Result: """ super().profile_pass() - # Reset optimization state for new optimization cycle - self.reset_optimization_state() - # Log profiling results if hasattr(self, "_initial_profiler_results") and self._initial_profiler_results: self.get_logger().record( - "profile_pass_complete", {"profiler_results": self._initial_profiler_results, "top_n": self.top_n} + "profile_pass_complete", + { + "profiler_results": self._initial_profiler_results, + "top_n": self.top_n, + }, ) def instrument_pass(self) -> Result: @@ -105,7 +309,11 @@ def instrument_pass(self) -> Result: # Log instrumentation completion self.get_logger().record( - "instrument_pass_complete", {"success": True, "note": "Instrumentation pass completed via parent class"} + "instrument_pass_complete", + { + "success": True, + "note": "Instrumentation pass completed via parent class", + }, ) return Result( @@ -125,12 +333,56 @@ def build_pass(self, validate_build_result=True) -> Result: Result: Build status and the output file path """ result = super().build(validate_build_result=validate_build_result) + + # Log build result if not result: logging.debug(f"Setting current summary to: {result.error_report}") self.current_summary = result.error_report + self.get_logger().record( + "build_pass_failed", + { + "success": False, + "error_report": result.error_report, + "kernel_files": getattr(self, "current_kernel_files", []), + }, + ) + # Add compilation failure to history so LLM learns from the error + if hasattr(self, "current_kernel_files") and self.current_kernel_files: + diff = self.compute_diff(self.current_kernel_files) + 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": 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')}", + optimized_code=failed_code, + ) + else: + self.get_logger().record( + "build_pass_success", + { + "success": True, + "kernel_files": getattr(self, "current_kernel_files", []), + }, + ) return result - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: + def optimize_pass( + self, + temperature: float = 0.0, + max_tokens: int = 3000, + target_kernel: str = None, + ) -> Result: """ Optimize the kernel to remove atomic contention via OpenAI API @@ -142,26 +394,14 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target Result: Optimized kernel as a file path """ super().optimize_pass() - llm_key = get_llm_api_key() system_prompt = ( - "You are a skilled GPU HIP programmer. Given a kernel," - " you will optimize it to remove atomic contention as much as possible" - " and provide a correct performant implementation. Do not modify" - " the kernel signature. Do not touch any other code, licenses, copyrights, or comments in the file." - " If you remove the copyright, your solution will be rejected." - " Do not include any markdown code blocks or text other than the code." + "You are a skilled GPU HIP programmer specializing in optimizing kernels " + "to reduce atomic contention and improve performance." ) - model = self.model - provider = self.provider - llm = LLM( - api_key=llm_key, - system_prompt=system_prompt, - model=model, - provider=provider, - logger=self.get_logger(), # Add logger here - ) + # Get LLM instance (initialized once per formula) + llm = self.get_llm(system_prompt) kernel = None kernel_file = None @@ -187,6 +427,12 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target kernel = filtered_report_card[0]["kernel"] self._parse_kernel_signature(kernel) + + # Store baseline metrics on first run + if self.baseline_atomic_latency is None: + self.baseline_atomic_latency = filtered_report_card[0]["atomics"]["atomic_lat"] + self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 + files = filtered_report_card[0]["source"]["files"] kernel_name = get_kernel_name(kernel) kernel_file = None @@ -204,32 +450,10 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target if kernel_file is None: return Result(success=False, error_report="Kernel file not found.") - user_prompt = ( - f"OPTIMIZATION ITERATION: You are optimizing kernel '{kernel}' to reduce atomic contention.\n\n" - f"SOURCE CODE:\n{unoptimized_file_content}\n\n" - ) - - # Add iteration context if this is not the first attempt - if self.current_summary is not None: - user_prompt += ( - f"PREVIOUS OPTIMIZATION RESULTS:\n{self.current_summary}\n\n" - "The previous optimization approach did not improve atomic contention. " - "Try a different strategy.\n\n" - ) - - # Add optimization history context - history_context = self.get_optimization_history_context() - user_prompt += f"{history_context}\n\n" - - # Add diff information to show what was changed - cur_diff = self.compute_diff([kernel_file]) - if cur_diff: - user_prompt += f"CHANGES FROM ORIGINAL CODE:\n{cur_diff}\n\n" - - user_prompt += ( - "OBJECTIVE: Reduce atomic contention while maintaining or improving runtime performance.\n" - "CONSTRAINTS: Preserve kernel signature and program semantics.\n" - "Return only the optimized code without any markdown or explanatory text." + # Build problem description + problem_description = ( + f"High atomic contention detected in kernel {kernel}. " + f"Please optimize to reduce atomic contention but do not change the semantics." ) self.previous_source_code = unoptimized_file_content @@ -257,34 +481,63 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target self.current_kernel_files = [kernel_file] - logging.debug(f"LLM prompt: {user_prompt}") logging.debug(f"System prompt: {system_prompt}") + logging.debug(f"Problem description: {problem_description}") try: - record_meta = f"Iteration {self.iteration_count + 1}" - optimized_file_content = llm.ask(user_prompt, record_meta=record_meta).strip() - optimized_file_content = self.postprocess_llm_code(optimized_file_content) - - # Track this optimization attempt - self.iteration_count += 1 + # Use DSPy signature with history - pass inputs as kwargs + # Build failure analysis from history + tracker_dict = self.optimization_tracker.to_dict() + failed_steps = [s for s in tracker_dict.get("steps", []) if not s.get("success", False)] + previous_failures_text = "" + if failed_steps: + previous_failures_text = "Previous failed attempts:\n" + for i, step in enumerate(failed_steps[-3:], 1): # Show last 3 failures + previous_failures_text += f"\nAttempt {i}: {step.get('report', 'Unknown error')}\n" + else: + previous_failures_text = "No previous failures yet. This is an early iteration." - # Store attempt info (will be updated with results after validation) - self.optimization_attempts.append( - { - "iteration": self.iteration_count, - "code": optimized_file_content, - "success": None, # Will be updated after performance validation - "improvement": 0.0, # Will be updated after performance validation - } + # Get latest diff if available + latest_diff_text = "" + if tracker_dict.get("steps"): + latest_step = tracker_dict["steps"][-1] + latest_diff_text = latest_step.get("diff", "No diff available") + else: + latest_diff_text = "No previous attempts yet." + + response = llm.ask( + signature=AtomicContentionOptimization, + answer_type=None, # Return full response object + record_meta="atomic_contention_optimization", + kernel_code=unoptimized_file_content, + code_preservation_rules="", # Field description contains the critical rules + output_format_rules="", # Field description contains the format requirements + problem_description=problem_description, + baseline_atomic_latency=str(self.baseline_atomic_latency), + baseline_time_ms=str(self.baseline_time_ms), + history=self.optimization_tracker.get_dspy_history(), + previous_failures=previous_failures_text, + latest_diff=latest_diff_text, + optimization_techniques="", # Field description contains the full content + amd_specific_optimizations="", # Field description contains the full content + optimization_categories="", # Field description contains the full content ) + # Extract optimized code from response + if hasattr(response, "optimized_code"): + optimized_file_content = response.optimized_code.strip() + else: + # Fallback for string response + optimized_file_content = str(response).strip() + + optimized_file_content = self.postprocess_llm_code(optimized_file_content) + # Log successful optimization self.get_logger().record( "optimization_success", { "optimized_code_length": len(optimized_file_content), "kernel_file": kernel_file, - "iteration": self.iteration_count, }, ) @@ -300,7 +553,10 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target ) except Exception as e: error_msg = f"An unexpected error occurred - {str(e)}" - self.get_logger().record("optimization_error", {"error": error_msg, "error_type": type(e).__name__}) + self.get_logger().record( + "optimization_error", + {"error": error_msg, "error_type": type(e).__name__}, + ) logging.error(error_msg) return Result(success=False, error_report=error_msg) @@ -323,9 +579,50 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) Result: Validation status """ result = super().correctness_validation_pass(self.current_kernel, self.current_args, accordo_absolute_tolerance) + + # Log correctness validation result if not result: logging.info(f"Setting current summary to: {result.error_report}") self.current_summary = result.error_report + self.get_logger().record( + "correctness_validation_failed", + { + "success": False, + "error_report": result.error_report, + "kernel": self.current_kernel, + "args": self.current_args, + }, + ) + # Add correctness failure to history so LLM learns from the error + if hasattr(self, "current_kernel_files") and self.current_kernel_files: + diff = self.compute_diff(self.current_kernel_files) + with open(self.current_kernel_files[0], "r") as f: + failed_code = f.read() + + 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_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')}", + optimized_code=failed_code, + ) + else: + self.get_logger().record( + "correctness_validation_success", + { + "success": True, + "kernel": self.current_kernel, + "args": self.current_args, + }, + ) return result def performance_validation_pass(self) -> Result: @@ -399,35 +696,7 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - if not success or speedup < 1: - self.current_summary = self.optimization_report - - # Update the latest optimization attempt with results - if self.optimization_attempts: - latest_attempt = self.optimization_attempts[-1] - latest_attempt["success"] = False - latest_attempt["improvement"] = -cycle_latency_improvement # Negative for worsening - logging.info( - f"Optimization iteration {latest_attempt['iteration']} FAILED: " - f"Atomic contention increased by {abs(cycle_latency_improvement):.1f}%" - ) - - return Result(success=False, error_report=self.optimization_report) - - # Update the latest optimization attempt with results - if self.optimization_attempts: - latest_attempt = self.optimization_attempts[-1] - latest_attempt["success"] = True - latest_attempt["improvement"] = cycle_latency_improvement - logging.info( - f"Optimization iteration {latest_attempt['iteration']} SUCCEEDED: " - f"Atomic contention reduced by {cycle_latency_improvement:.1f}%" - ) - - # Log performance validation results - optimized_code_string = None - if self.optimization_attempts: - optimized_code_string = self.optimization_attempts[-1].get("code") + # Log performance validation results (always, even if failed) self.get_logger().record( "performance_validation_complete", { @@ -440,22 +709,88 @@ def performance_validation_pass(self) -> Result: "metric_improvement": metric_improvement, "cycle_latency_improvement": cycle_latency_improvement, "optimization_report": self.optimization_report, - "optimized_code_string": optimized_code_string, }, ) logging.info(self.optimization_report) - self.success = True - return Result(success=True, asset={"log": 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 + return Result(success=False, error_report=self.optimization_report) def write_results(self, output_file: str = None): """ - Writes the results to the output file. + Writes the results to the output file using the best optimization attempt. """ + # Restore best results for output + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report + + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) + + # 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": 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.success}, + additional_results={ + "formula": "atomicContention", + "success": self.success, + "optimization_history": self.optimization_tracker.to_dict(), + **metric_fields, + }, ) def summarize_previous_passes(self): @@ -463,34 +798,3 @@ def summarize_previous_passes(self): Summarizes the results of the previous passes for future prompts. """ pass - - def get_optimization_history_context(self) -> str: - """ - Generate context about previous optimization attempts for the LLM. - - Returns: - str: Formatted context about previous optimization attempts - """ - if not self.optimization_attempts: - return "This is the first optimization attempt." - - context = "PREVIOUS OPTIMIZATION ATTEMPTS:\n" - for i, attempt in enumerate(self.optimization_attempts, 1): - context += f"Iteration {i}: " - if attempt["success"]: - context += f"SUCCESS - Atomic contention reduced by {attempt['improvement']:.1f}%\n" - else: - context += f"FAILED - Atomic contention increased by {abs(attempt['improvement']):.1f}%\n" - - return context - - def reset_optimization_state(self): - """ - Reset the optimization state for a new optimization cycle. - """ - self.optimization_attempts = [] - self.iteration_count = 0 - self.current_summary = None - self.previous_source_code = None - self.success = False - logging.info("Optimization state reset for new cycle") diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index 95648491..b096d768 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -28,18 +28,209 @@ import os import shutil -from intelliperf.core.llm import LLM +import dspy + from intelliperf.formulas.formula_base import ( Formula_Base, + OptimizationTracker, Result, filter_json_field, get_kernel_name, ) -from intelliperf.utils.env import get_llm_api_key from intelliperf.utils.process import capture_subprocess_output from intelliperf.utils.regex import generate_ecma_regex_from_list +class BankConflictOptimization(dspy.Signature): + """Optimize GPU kernel code to reduce shared memory bank conflicts and improve performance.""" + + kernel_code = dspy.InputField(desc="The current kernel source code that needs bank conflict optimization.") + + code_preservation_rules = dspy.InputField( + desc="""CRITICAL: These rules MUST be followed to avoid compilation errors: + +1. **NEVER modify, abbreviate, or remove copyright/license headers** + - Copy the ENTIRE copyright header exactly as written, character by character + - DO NOT use '...' or any ellipsis inside comments - this creates unterminated comment errors + - If you see: /* Copyright (c) 2025... */ you MUST write the full text + +2. **NEVER use '...' or ellipsis in C/C++ comments** + - The sequence '...' does NOT close a /* */ comment block + - This causes: error: unterminated /* comment + - Either copy the full comment text or remove it entirely (but keep copyright!) + +3. **Preserve the EXACT kernel signature** + - Do not change: function name, parameter types, parameter names, or return type + - The kernel must be callable with the same arguments as before""" + ) + + output_format_rules = dspy.InputField( + desc="""CRITICAL: Output format requirements: + +1. **No markdown code blocks or formatting** + - Do NOT include: ```cpp, ```c, ```hip, or ``` markers + - Do NOT include any explanatory text before or after the code + +2. **Output ONLY the complete, compilable source code** + - Include all necessary #include statements + - Include the full copyright header (copied verbatim) + - Include all function implementations + - The output must be ready to write directly to a .hip file + +3. **Preserve all existing comments and licenses** + - Copy them exactly as they appear in the original code""" + ) + + problem_description = dspy.InputField( + desc="Description of the bank conflict issue detected, including specific shared memory access patterns causing conflicts and performance impact." + ) + + # Baseline metrics for reference + baseline_bank_conflicts = dspy.InputField( + desc="Baseline bank conflict ratio. Indicates how many shared memory accesses result in bank conflicts - higher values mean more conflicts and worse performance. This is the initial unoptimized value before any optimization attempts." + ) + baseline_time_ms = dspy.InputField( + desc="Baseline kernel execution time in milliseconds. This is the initial unoptimized kernel runtime before any optimization attempts." + ) + + # History of all previous attempts + history = dspy.History = dspy.InputField( + desc="Complete history of previous optimization attempts, including the code changes (diffs), before/after bank conflict ratios, speedup ratios, and whether each attempt improved or regressed performance. Use this to avoid repeating failed approaches and build on successful patterns." + ) + + previous_failures = dspy.InputField( + desc="CRITICAL: Analysis of previous failed attempts. DO NOT repeat these mistakes: compilation errors (unterminated comments, missing semicolons), correctness failures (changed semantics, wrong outputs), or performance regressions (same approach with minor variations). Learn from these failures and try fundamentally different optimization strategies." + ) + + latest_diff = dspy.InputField( + desc="The most recent code change that was attempted. If this resulted in a failure, you MUST try a completely different approach. Do not make minor variations of the same optimization - innovate with a new strategy." + ) + + # Low-level optimization techniques + optimization_techniques = dspy.InputField( + desc="""Available low-level bank conflict optimization techniques to consider: + +1. **Non-temporal loads**: Treat read-once data as streaming, avoid cache pollution. + Example: float x = __builtin_nontemporal_load(&input[i]); + +2. **Warp-wide cooperative loads** using shuffles: + float val = __shfl_sync(0xffffffff, local_data, 0); + +3. **Prefetching into registers**: + float prefetch = input[i + 1]; // Use later + +4. **Manual register tiling with unrolled loops**: + float tile[4]; + #pragma unroll + for (int i = 0; i < 4; ++i) tile[i] = input[idx + i]; + +5. **Double-buffering** for overlapping memory + compute: + load_tile(buf_a); + load_tile(buf_b); compute(buf_a); swap(); + +6. **Structure-of-Arrays transformation** to align thread memory access: + float val = soa.x[threadIdx.x]; // Coalesced + +7. **Aligned vectorized memory access** using float4: + float4 x = reinterpret_cast(input)[idx];""" + ) + + amd_specific_optimizations = dspy.InputField( + desc="""AMD-specific optimizations for CDNA GPUs (MI250X, MI300X): + +**AMD MFMA (Matrix Fused Multiply-Add) intrinsics**: +These instructions execute on a **wavefront-wide** basis (64 threads), using per-lane vector registers to load parts of matrices A, B, C. + +**MFMA Intrinsic Syntax:** + d = __builtin_amdgcn_mfma__(a, b, c, cbsz, abid, blgp); + +**Parameters:** +- CDfmt: format of C and D (e.g., f32, fp32) +- ABfmt: format of A and B (e.g., f16, bf16, i8) +- M, N, K: matrix tile dimensions +- a, b, c: registers or scalars from matrices A, B, C +- d: output register of resulting matrix tile D +- cbsz: broadcast size (0=no broadcast, 1=2-wide broadcast, etc.) +- abid: A-matrix block to broadcast from (used with cbsz) +- blgp: B-matrix swizzle pattern (0=normal, 1=lanes 0-31β†’32-63, 2=lanes 32-63β†’0-31, 3=rotate down by 16, 4-7=group-wide broadcast) + +**Example (FP32, 16x16x4 GEMM):** +```cpp +__global__ void sgemm_16x16x4(const float *A, const float *B, float *D) { + using float4 = __attribute__((__vector_size__(4 * sizeof(float)))) float; + float4 d = {0}; + + int mk = threadIdx.y + 4 * threadIdx.x; + int kn = threadIdx.x + 16 * threadIdx.y; + + float amk = A[mk]; + float bkn = B[kn]; + d = __builtin_amdgcn_mfma_f32_16x16x4f32(amk, bkn, d, 0, 0, 0); + + for (int i = 0; i < 4; ++i) + D[threadIdx.x + i * 16 + threadIdx.y * 4 * 16] = d[i]; +} +``` +Launch with: dim3 block(16, 4); dim3 grid(1, 1); + +**Performance (CDNA2/CDNA3 - MI250X/MI300X):** +- FP32 16x16x4: 256 flops/cycle/CU +- FP16/BF16 16x16x16: 1024 flops/cycle/CU +- INT8 16x16x16: 1024 flops/cycle/CU +- FP64 4x4x4: 128-256 flops/cycle/CU + +**Key Benefits:** +- MFMA offers 2-4Γ— throughput over vector FMAs +- Use wavefront-aligned loads and register blocking to match instruction layout +- Avoid using warp-level primitives with the _sync suffix (e.g., __shfl_sync, __ballot_sync) and use the non-sync variants instead for broader compatibility +- We are optimizing for CDNA GPUs (MI300X) so target this architecture + +**LDS Bank Conflict Avoidance on AMD:** +- AMD GCN/CDNA architectures have 32 LDS banks (vs NVIDIA's 32) +- Each bank is 4 bytes wide +- Conflicts occur when multiple threads in a wavefront access different addresses in the same bank +- Use padding or stride adjustments to avoid bank conflicts""" + ) + + optimization_categories = dspy.InputField( + desc="""Categories of optimization strategies to explore: + +1. **Bank Conflict Resolution Strategies:** + - Padding shared memory arrays to avoid conflicts + - Changing access strides to distribute across banks + - Transposing data layout in shared memory + - Using conflict-free indexing schemes + - Reorganizing thread-to-data mappings + +2. **Memory Access Pattern Changes:** + - Non-temporal loads for streaming data (__ldg, __builtin_nontemporal_load) + - Warp-wide cooperative loads with shuffles (__shfl, __shfl_down) + - Manual register tiling and double-buffering + - Structure-of-Arrays transformations + +3. **Advanced Vectorization:** + - Vectorized memory access patterns (float4, int4) + - SIMD-friendly loop unrolling + - Memory coalescing with vector types + +4. **AMD-Specific Optimizations:** + - AMD MFMA (Matrix Fused Multiply-Add) intrinsics for CDNA GPUs (MI250X, MI300X) + - Wavefront-level optimizations (64-thread blocks) + - AMD-specific LDS bank organization (32 banks, 4 bytes each) + - AMD-specific memory hierarchy optimizations + +5. **Algorithmic Changes:** + - Different loop ordering strategies + - Alternative memory layout schemes + - New data structure organizations + - Wavefront/warp specialization and/or workgroup/block specialization""" + ) + + optimized_code = dspy.OutputField( + desc="The complete, runnable optimized kernel code with reduced bank conflicts. Must preserve the exact kernel signature, all comments, licenses, and copyright notices. The code should eliminate or minimize shared memory bank conflicts through better access patterns or padding." + ) + + class bank_conflict(Formula_Base): def __init__( self, @@ -78,6 +269,28 @@ def __init__( self.previous_source_code = None self.success = False + # Initialize optimization tracker + # Bank conflict optimization maximizes conflict reduction (minimize bank conflicts) + # Automatically calculates conflict_improvement from unoptimized_conflicts / optimized_conflicts + self.optimization_tracker = OptimizationTracker( + max_iterations=self.num_attempts, + primary_metric="conflict_improvement", + maximize=True, + before_metric="unoptimized_conflicts", + after_metric="optimized_conflicts", + ) + + # Store baseline metrics (set during first profiling) + 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. @@ -89,8 +302,47 @@ def build_pass(self, validate_build_result=True) -> Result: Result: Build status and the output file path """ result = super().build(validate_build_result=validate_build_result) + + # Log build result if not result: self.current_summary = result.error_report + self.get_logger().record( + "build_pass_failed", + { + "success": False, + "error_report": result.error_report, + "kernel_files": getattr(self, "current_kernel_files", []), + }, + ) + # Add compilation failure to history so LLM learns from the error + if hasattr(self, "current_kernel_files") and self.current_kernel_files: + diff = self.compute_diff(self.current_kernel_files) + 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": 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", + { + "success": True, + "kernel_files": getattr(self, "current_kernel_files", []), + }, + ) return result def profile_pass(self) -> Result: @@ -105,7 +357,11 @@ def profile_pass(self) -> Result: # Log profiling results if hasattr(self, "_initial_profiler_results") and self._initial_profiler_results: self.get_logger().record( - "profile_pass_complete", {"profiler_results": self._initial_profiler_results, "top_n": self.top_n} + "profile_pass_complete", + { + "profiler_results": self._initial_profiler_results, + "top_n": self.top_n, + }, ) def get_top_kernel(self) -> str: @@ -133,7 +389,11 @@ def instrument_pass(self) -> Result: # Log instrumentation completion self.get_logger().record( - "instrument_pass_complete", {"success": True, "note": "Instrumentation pass completed via parent class"} + "instrument_pass_complete", + { + "success": True, + "note": "Instrumentation pass completed via parent class", + }, ) return Result( @@ -201,7 +461,12 @@ def instrument_pass(self) -> Result: ) return Result(success=True, asset=self._instrumentation_results) - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: + def optimize_pass( + self, + temperature: float = 0.0, + max_tokens: int = 3000, + target_kernel: str = None, + ) -> Result: """ Optimize the kernel to remove shared memory bank conflicts via OpenAI API @@ -213,22 +478,14 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target Result: Optimized kernel as a file path """ super().optimize_pass() - llm_key = get_llm_api_key() system_prompt = ( - "You are a skilled GPU HIP programmer. Given a kernel," - " you will optimize it to remove shared memory bank conflicts" - " and provide a correct performant implementation. Do not modify" - " the kernel signature. Do not touch any other code, licenses, copyrights, or comments in the file." - " If you remove the copyright, your solution will be rejected." - " Do not include any markdown code blocks or text other than the code." + "You are a skilled GPU HIP programmer specializing in optimizing kernels " + "to reduce shared memory bank conflicts and improve performance." ) - provider = self.provider - model = self.model - llm = LLM( - api_key=llm_key, system_prompt=system_prompt, model=model, provider=provider, logger=self.get_logger() - ) + # Get LLM instance (initialized once per formula) + llm = self.get_llm(system_prompt) kernel_to_optimize = self.get_top_kernel() if kernel_to_optimize is None: @@ -258,6 +515,12 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target kernel = filtered_report_card[0]["kernel"] self._parse_kernel_signature(kernel) + + # Store baseline metrics on first run + if self.baseline_bank_conflicts is None: + self.baseline_bank_conflicts = filtered_report_card[0]["lds"]["bc"] + self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 + files = filtered_report_card[0]["source"]["files"] kernel_name = get_kernel_name(kernel) kernel_file = None @@ -276,16 +539,11 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target if kernel_file is None: return Result(success=False, error_report="Kernel file not found.") - user_prompt = ( - f"There is a bank conflict in the kernel {kernel} in the source code {unoptimized_file_content}." - f" Please fix the conflict but do not change the semantics of the program." - " Do not remove any comments or licenses." - " Do not include any markdown code blocks or text other than the code." + # Build problem description + problem_description = ( + f"Shared memory bank conflicts detected in kernel {kernel}. " + f"Please optimize to reduce bank conflicts but do not change the semantics." ) - if self.current_summary is not None: - user_prompt += f"\n\nThe current summary is: {self.current_summary}" - cur_diff = self.compute_diff([kernel_file]) - user_prompt += f"\nThe diff between the current and initial code is: {cur_diff}" self.previous_source_code = unoptimized_file_content @@ -312,16 +570,63 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target self.current_kernel_files = [kernel_file] logging.debug(f"System prompt: {system_prompt}") - logging.debug(f"LLM prompt: {user_prompt}") + logging.debug(f"Problem description: {problem_description}") try: - optimized_file_content = llm.ask(user_prompt).strip() + # Use DSPy signature with history - pass inputs as kwargs + # Build failure analysis from history + tracker_dict = self.optimization_tracker.to_dict() + failed_steps = [s for s in tracker_dict.get("steps", []) if not s.get("success", False)] + previous_failures_text = "" + if failed_steps: + previous_failures_text = "Previous failed attempts:\n" + for i, step in enumerate(failed_steps[-3:], 1): # Show last 3 failures + previous_failures_text += f"\nAttempt {i}: {step.get('report', 'Unknown error')}\n" + else: + previous_failures_text = "No previous failures yet. This is an early iteration." + + # Get latest diff if available + latest_diff_text = "" + if tracker_dict.get("steps"): + latest_step = tracker_dict["steps"][-1] + latest_diff_text = latest_step.get("diff", "No diff available") + else: + latest_diff_text = "No previous attempts yet." + + response = llm.ask( + signature=BankConflictOptimization, + answer_type=None, # Return full response object + record_meta="bank_conflict_optimization", + kernel_code=unoptimized_file_content, + code_preservation_rules="", # Field description contains the critical rules + output_format_rules="", # Field description contains the format requirements + problem_description=problem_description, + baseline_bank_conflicts=str(self.baseline_bank_conflicts), + baseline_time_ms=str(self.baseline_time_ms), + history=self.optimization_tracker.get_dspy_history(), + previous_failures=previous_failures_text, + latest_diff=latest_diff_text, + optimization_techniques="", # Field description contains the full content + amd_specific_optimizations="", # Field description contains the full content + optimization_categories="", # Field description contains the full content + ) + + # Extract optimized code from response + if hasattr(response, "optimized_code"): + optimized_file_content = response.optimized_code.strip() + else: + # Fallback for string response + optimized_file_content = str(response).strip() + optimized_file_content = self.postprocess_llm_code(optimized_file_content) # Log successful optimization self.get_logger().record( "optimization_success", - {"optimized_code_length": len(optimized_file_content), "kernel_file": kernel_file}, + { + "optimized_code_length": len(optimized_file_content), + "kernel_file": kernel_file, + }, ) with open(kernel_file, "w") as f: @@ -336,7 +641,10 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target ) except Exception as e: error_msg = f"An unexpected error occurred - {str(e)}" - self.get_logger().record("optimization_error", {"error": error_msg, "error_type": type(e).__name__}) + self.get_logger().record( + "optimization_error", + {"error": error_msg, "error_type": type(e).__name__}, + ) logging.error(error_msg) return Result(success=False, error_report=error_msg) @@ -359,8 +667,49 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) Result: Validation status """ result = super().correctness_validation_pass(self.current_kernel, self.current_args, accordo_absolute_tolerance) + + # Log correctness validation result if not result: self.current_summary = result.error_report + self.get_logger().record( + "correctness_validation_failed", + { + "success": False, + "error_report": result.error_report, + "kernel": self.current_kernel, + "args": self.current_args, + }, + ) + # Add correctness failure to history so LLM learns from the error + if hasattr(self, "current_kernel_files") and self.current_kernel_files: + diff = self.compute_diff(self.current_kernel_files) + with open(self.current_kernel_files[0], "r") as f: + failed_code = f.read() + + 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", + { + "success": True, + "kernel": self.current_kernel, + "args": self.current_args, + }, + ) return result def performance_validation_pass(self) -> Result: @@ -387,6 +736,7 @@ def performance_validation_pass(self) -> Result: 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 if unoptimized_conflicts != 0 else 0 ) * 100 @@ -423,7 +773,7 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Log performance validation results + # Log performance validation results (always, even if failed) self.get_logger().record( "performance_validation_complete", { @@ -433,27 +783,91 @@ def performance_validation_pass(self) -> Result: "unoptimized_conflicts": unoptimized_conflicts, "optimized_conflicts": optimized_conflicts, "speedup": speedup, + "conflict_improvement": conflict_improvement, "conflict_improvement_percentage": conflict_improvement_percentage, "optimization_report": self.optimization_report, }, ) - if not success or speedup < 1: - self.current_summary = self.optimization_report - return Result(success=False, error_report=self.optimization_report) - logging.info(self.optimization_report) - self.success = True - return Result(success=True, asset={"log": 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 + return Result(success=False, error_report=self.optimization_report) def write_results(self, output_file: str = None): """ - Writes the results to the output file. + Writes the results to the output file using the best optimization attempt. """ + # Restore best results for output + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report + + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) + + # 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": 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.success}, + additional_results={ + "formula": "bankConflict", + "success": self.success, + "optimization_history": self.optimization_tracker.to_dict(), + **metric_fields, + }, ) def summarize_previous_passes(self): diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index b836f4fa..7fcc053c 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -29,7 +29,9 @@ import sys import time from abc import abstractmethod +from dataclasses import asdict, dataclass, field from pprint import pformat +from typing import List, Optional import ml_dtypes import numpy as np @@ -44,6 +46,143 @@ from intelliperf.utils.process import capture_subprocess_output, exit_on_fail +@dataclass +class OptimizationStep: + """Represents a single optimization attempt""" + + iteration: int + diff: str + report: str + metrics: dict + success: bool + timestamp: float = field(default_factory=time.time) + + def get_metric(self, key: str, default=0.0): + """Helper to safely get metrics""" + return self.metrics.get(key, default) + + +class OptimizationTracker: + """Tracks optimization history using dspy.History for proper conversation management""" + + def __init__( + self, + max_iterations: int = 10, + primary_metric: str = "speedup", + maximize: bool = True, + before_metric: Optional[str] = None, + after_metric: Optional[str] = None, + ): + self.steps: List[OptimizationStep] = [] + self.current_iteration: int = 0 + self.max_iterations: int = max_iterations + self.best_step: Optional[OptimizationStep] = None + self.primary_metric: str = primary_metric + self.maximize: bool = maximize + self.initial_source_code: Optional[str] = None + # For auto-calculating improvements from before/after metrics + self.before_metric: Optional[str] = before_metric + self.after_metric: Optional[str] = after_metric + + # History messages for DSPy (stored as list of dicts) + self.history_messages = [] + + def add_step( + self, + diff: str, + report: str, + metrics: dict, + success: bool, + request: str = "", + optimized_code: str = "", + ) -> OptimizationStep: + """Add step and auto-update best based on primary metric""" + # Auto-calculate improvement if before/after metrics are configured + if self.before_metric and self.after_metric: + before = metrics.get(self.before_metric, 0) + after = metrics.get(self.after_metric, 0) + if before != 0: + improvement = after / before if after != 0 else 1.0 + metrics[self.primary_metric] = improvement + + step = OptimizationStep( + iteration=self.current_iteration, + diff=diff, + report=report, + metrics=metrics, + success=success, + ) + self.steps.append(step) + self.current_iteration += 1 + + # Add to DSPy history with explicit counters for better LLM learning + # Extract common metrics for structured history + history_entry = { + "iteration": self.current_iteration - 1, + "request": request, + "optimized_code": optimized_code, + "result_summary": report, + "diff": diff, + "success": "βœ“ Improved" if success else "βœ— Regressed", + } + + # Add explicit before/after counters if available + if self.before_metric and self.after_metric: + before_val = metrics.get(self.before_metric, 0) + after_val = metrics.get(self.after_metric, 0) + improvement_val = metrics.get(self.primary_metric, 1.0) + + history_entry.update( + { + f"before_{self.before_metric}": before_val, + f"after_{self.after_metric}": after_val, + f"{self.primary_metric}": improvement_val, + } + ) + + # Add all other metrics + history_entry["all_metrics"] = metrics + + self.history_messages.append(history_entry) + + # Auto-update best + if self.best_step is None: + self.best_step = step + else: + new_val = step.get_metric(self.primary_metric) + cur_val = self.best_step.get_metric(self.primary_metric) + + if (self.maximize and new_val > cur_val) or (not self.maximize and new_val < cur_val): + self.best_step = step + + return step + + def get_dspy_history(self): + """Get the DSPy history messages for use in signatures""" + + # Return a simple object with messages attribute for DSPy + class HistoryWrapper: + def __init__(self, messages): + self.messages = messages + + return HistoryWrapper(self.history_messages) + + def has_reached_max_iterations(self) -> bool: + """Check if max iterations reached""" + return self.current_iteration >= self.max_iterations + + def to_dict(self) -> dict: + """Serialize for JSON output""" + return { + "steps": [asdict(step) for step in self.steps], + "best_step": asdict(self.best_step) if self.best_step else None, + "current_iteration": self.current_iteration, + "max_iterations": self.max_iterations, + "primary_metric": self.primary_metric, + "maximize": self.maximize, + } + + class Result: def __init__(self, success: bool, error_report: str = "", asset=None): self.success: bool = success @@ -84,6 +223,7 @@ def __init__( provider: str = "openai", in_place: bool = False, unittest_command: str = None, + num_attempts: int = 10, ): # Private self.__name = name # name of the run @@ -93,6 +233,9 @@ def __init__( logging.debug(f"project_directory: {project_directory}") logging.debug(f"app_cmd: {app_cmd}") + # Store num_attempts + self.num_attempts = num_attempts + # Initialize logger self._logger = Logger(run_name=name) self._logger.record( @@ -107,6 +250,7 @@ def __init__( "model": model, "provider": provider, "in_place": in_place, + "num_attempts": num_attempts, }, ) @@ -144,6 +288,10 @@ def __init__( self.current_args = None self.current_kernel_signature = None + # Initialize DSPy LLM once per formula (lazily, only if needed) + self._llm = None + self._dspy_configured = False + self.build() def get_logger(self) -> Logger: @@ -151,10 +299,37 @@ def get_logger(self) -> Logger: Get the logger instance for this formula run. Returns: - Logger: The logger instance + Logger: The logger instance """ return self._logger + def get_llm(self, system_prompt: str): + """ + Get or create the LLM instance for this formula (lazy initialization). + + DSPy is configured once per formula instance for efficiency. + + Args: + system_prompt: System prompt for the LLM + + Returns: + LLM: The LLM instance configured for this formula + """ + from intelliperf.core.llm import LLM + from intelliperf.utils.env import get_llm_api_key + + if self._llm is None: + llm_key = get_llm_api_key() + self._llm = LLM( + api_key=llm_key, + system_prompt=system_prompt, + model=self.model, + provider=self.provider, + logger=self.get_logger(), + ) + logging.debug(f"Initialized LLM once for formula: {self.model} via {self.provider}") + return self._llm + def _parse_kernel_signature(self, kernel_signature: str): """ Parses a kernel signature to extract the kernel name and its arguments. @@ -394,10 +569,10 @@ def postprocess_llm_code(self, optimized_file_content: str) -> str: Removes markdown code blocks (```c++, ```python, etc.) from the LLM response. Args: - optimized_file_content (str): The LLM generated code + optimized_file_content (str): The LLM generated code Returns: - str: The post-processed code + str: The post-processed code """ # Remove markdown code blocks if present content = optimized_file_content.strip() @@ -456,7 +631,12 @@ def inplace_update(self, filepaths: list[str]): with open(reference_filepath, "w") as f: f.write(optimized_content) - def write_results(self, output_file: str = None, additional_results: dict = {}, diagnose_only: bool = False): + def write_results( + self, + output_file: str = None, + additional_results: dict = {}, + diagnose_only: bool = False, + ): """ Writes the results to the output file. """ @@ -480,6 +660,37 @@ def write_results(self, output_file: str = None, additional_results: dict = {}, write_results(results, output_file) +def _add_diff_lines_recursive(obj): + """ + Recursively add diff_lines field to any dict that contains a diff field. + + Args: + obj: Object to process (dict, list, or other) + + Returns: + Processed object with diff_lines added where applicable + """ + if isinstance(obj, dict): + # Process all values in the dict recursively + result = {} + for key, value in obj.items(): + result[key] = _add_diff_lines_recursive(value) + + # Add diff_lines if diff exists and diff_lines doesn't + if "diff" in result and "diff_lines" not in result: + diff_value = result["diff"] + if isinstance(diff_value, str): + result["diff_lines"] = diff_value.split("\n") if diff_value else [] + + return result + elif isinstance(obj, list): + # Process all items in the list recursively + return [_add_diff_lines_recursive(item) for item in obj] + else: + # Return other types as-is + return obj + + def write_results(json_results: dict, output_file: str = None): """ Writes the results to the output file. @@ -487,6 +698,9 @@ def write_results(json_results: dict, output_file: str = None): log_message = f"Writing results to {output_file}" if output_file is not None else "Writing results to stdout" logging.info(log_message) + # Add diff_lines to all diffs in the results recursively + json_results = _add_diff_lines_recursive(json_results) + if output_file is None: print(json.dumps(json_results, indent=2)) elif output_file.endswith(".json"): diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index 9e57628b..3d5c2000 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -25,15 +25,193 @@ import json import logging import os +import sys + +import dspy -from intelliperf.core.llm import LLM from intelliperf.formulas.formula_base import ( Formula_Base, + OptimizationTracker, Result, filter_json_field, get_kernel_name, ) -from intelliperf.utils.env import get_llm_api_key + + +class MemoryAccessOptimization(dspy.Signature): + """Optimize GPU kernel code to improve memory coalescing and access patterns.""" + + kernel_code = dspy.InputField(desc="The current kernel source code that needs memory access optimization.") + + code_preservation_rules = dspy.InputField( + desc="""CRITICAL: These rules MUST be followed to avoid compilation errors: + +1. **NEVER modify, abbreviate, or remove copyright/license headers** + - Copy the ENTIRE copyright header exactly as written, character by character + - DO NOT use '...' or any ellipsis inside comments - this creates unterminated comment errors + - If you see: /* Copyright (c) 2025... */ you MUST write the full text + +2. **NEVER use '...' or ellipsis in C/C++ comments** + - The sequence '...' does NOT close a /* */ comment block + - This causes: error: unterminated /* comment + - Either copy the full comment text or remove it entirely (but keep copyright!) + +3. **Preserve the EXACT kernel signature** + - Do not change: function name, parameter types, parameter names, or return type + - The kernel must be callable with the same arguments as before""" + ) + + output_format_rules = dspy.InputField( + desc="""CRITICAL: Output format requirements: + +1. **No markdown code blocks or formatting** + - Do NOT include: ```cpp, ```c, ```hip, or ``` markers + - Do NOT include any explanatory text before or after the code + +2. **Output ONLY the complete, compilable source code** + - Include all necessary #include statements + - Include the full copyright header (copied verbatim) + - Include all function implementations + - The output must be ready to write directly to a .hip file + +3. **Preserve all existing comments and licenses** + - Copy them exactly as they appear in the original code""" + ) + + problem_description = dspy.InputField( + desc="Description of the memory access issue detected, including specific uncoalesced access patterns and performance impact." + ) + + # Baseline metrics for reference + baseline_coalesced_pct = dspy.InputField( + desc="Baseline memory coalescing efficiency (25-100%). Indicates how well memory instructions were coalesced by the address processing unit, ranging from uncoalesced (25%) to fully coalesced (100%). Calculated as the average number of thread-requests generated per instruction divided by the ideal number of thread-requests per instruction. This is the initial unoptimized value before any optimization attempts." + ) + baseline_time_ms = dspy.InputField( + desc="Baseline kernel execution time in milliseconds. This is the initial unoptimized kernel runtime before any optimization attempts." + ) + + # History of all previous attempts + history = dspy.History = dspy.InputField( + desc="Complete history of previous optimization attempts, including the code changes (diffs), before/after coalescing percentages, speedup ratios, and whether each attempt improved or regressed performance. Use this to avoid repeating failed approaches and build on successful patterns." + ) + + previous_failures = dspy.InputField( + desc="CRITICAL: Analysis of previous failed attempts. DO NOT repeat these mistakes: compilation errors (unterminated comments, missing semicolons), correctness failures (changed semantics, wrong outputs), or performance regressions (same approach with minor variations). Learn from these failures and try fundamentally different optimization strategies." + ) + + latest_diff = dspy.InputField( + desc="The most recent code change that was attempted. If this resulted in a failure, you MUST try a completely different approach. Do not make minor variations of the same optimization - innovate with a new strategy." + ) + + # Low-level optimization techniques + optimization_techniques = dspy.InputField( + desc="""Available low-level memory access optimization techniques to consider: + +1. **Non-temporal loads**: Treat read-once data as streaming, avoid cache pollution. + Example: float x = __builtin_nontemporal_load(&input[i]); + +2. **Warp-wide cooperative loads** using shuffles: + float val = __shfl_sync(0xffffffff, local_data, 0); + +3. **Prefetching into registers**: + float prefetch = input[i + 1]; // Use later + +4. **Manual register tiling with unrolled loops**: + float tile[4]; + #pragma unroll + for (int i = 0; i < 4; ++i) tile[i] = input[idx + i]; + +5. **Double-buffering** for overlapping memory + compute: + load_tile(buf_a); + load_tile(buf_b); compute(buf_a); swap(); + +6. **Structure-of-Arrays transformation** to align thread memory access: + float val = soa.x[threadIdx.x]; // Coalesced + +7. **Aligned vectorized memory access** using float4: + float4 x = reinterpret_cast(input)[idx];""" + ) + + amd_specific_optimizations = dspy.InputField( + desc="""AMD-specific optimizations for CDNA GPUs (MI250X, MI300X): + +**AMD MFMA (Matrix Fused Multiply-Add) intrinsics**: +These instructions execute on a **wavefront-wide** basis (64 threads), using per-lane vector registers to load parts of matrices A, B, C. + +**MFMA Intrinsic Syntax:** + d = __builtin_amdgcn_mfma__(a, b, c, cbsz, abid, blgp); + +**Parameters:** +- CDfmt: format of C and D (e.g., f32, fp32) +- ABfmt: format of A and B (e.g., f16, bf16, i8) +- M, N, K: matrix tile dimensions +- a, b, c: registers or scalars from matrices A, B, C +- d: output register of resulting matrix tile D +- cbsz: broadcast size (0=no broadcast, 1=2-wide broadcast, etc.) +- abid: A-matrix block to broadcast from (used with cbsz) +- blgp: B-matrix swizzle pattern (0=normal, 1=lanes 0-31β†’32-63, 2=lanes 32-63β†’0-31, 3=rotate down by 16, 4-7=group-wide broadcast) + +**Example (FP32, 16x16x4 GEMM):** +```cpp +__global__ void sgemm_16x16x4(const float *A, const float *B, float *D) { + using float4 = __attribute__((__vector_size__(4 * sizeof(float)))) float; + float4 d = {0}; + + int mk = threadIdx.y + 4 * threadIdx.x; + int kn = threadIdx.x + 16 * threadIdx.y; + + float amk = A[mk]; + float bkn = B[kn]; + d = __builtin_amdgcn_mfma_f32_16x16x4f32(amk, bkn, d, 0, 0, 0); + + for (int i = 0; i < 4; ++i) + D[threadIdx.x + i * 16 + threadIdx.y * 4 * 16] = d[i]; +} +``` +Launch with: dim3 block(16, 4); dim3 grid(1, 1); + +**Performance (CDNA2/CDNA3 - MI250X/MI300X):** +- FP32 16x16x4: 256 flops/cycle/CU +- FP16/BF16 16x16x16: 1024 flops/cycle/CU +- INT8 16x16x16: 1024 flops/cycle/CU +- FP64 4x4x4: 128-256 flops/cycle/CU + +**Key Benefits:** +- MFMA offers 2-4Γ— throughput over vector FMAs +- Use wavefront-aligned loads and register blocking to match instruction layout +- Avoid using warp-level primitives with the _sync suffix (e.g., __shfl_sync, __ballot_sync) and use the non-sync variants instead for broader compatibility +- We are optimizing for CDNA GPUs (MI300X) so target this architecture""" + ) + + optimization_categories = dspy.InputField( + desc="""Categories of optimization strategies to explore: + +1. **Memory Access Pattern Changes:** + - Non-temporal loads for streaming data (__ldg, __builtin_nontemporal_load) + - Warp-wide cooperative loads with shuffles (__shfl, __shfl_down) + - Manual register tiling and double-buffering + - Structure-of-Arrays transformations + +2. **Advanced Vectorization:** + - Vectorized memory access patterns (float4, int4) + - SIMD-friendly loop unrolling + - Memory coalescing with vector types + +3. **AMD-Specific Optimizations:** + - AMD MFMA (Matrix Fused Multiply-Add) intrinsics for CDNA GPUs (MI250X, MI300X) + - Wavefront-level optimizations (64-thread blocks) + - AMD-specific memory hierarchy optimizations + +4. **Algorithmic Changes:** + - Different loop ordering strategies + - Alternative memory layout schemes + - New data structure organizations + - Wavefront/warp specialization and/or workgroup/block specialization""" + ) + + optimized_code = dspy.OutputField( + desc="The complete, runnable optimized kernel code with improved memory access patterns. Must preserve the exact kernel signature, all comments, licenses, and copyright notices. The code should improve memory coalescing by ensuring threads in a warp access contiguous memory addresses." + ) class memory_access(Formula_Base): @@ -76,6 +254,28 @@ def __init__( self.previous_source_code = None self.success = False + # Initialize optimization tracker + # 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, + before_metric="unoptimized_coal", + after_metric="optimized_coal", + ) + + # Store baseline metrics (set during first profiling) + 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. @@ -87,8 +287,47 @@ def build_pass(self, validate_build_result=True) -> Result: Result: Build status and the output file path """ result = super().build(validate_build_result=validate_build_result) + + # Log build result if not result: self.current_summary = result.error_report + self.get_logger().record( + "build_pass_failed", + { + "success": False, + "error_report": result.error_report, + "kernel_files": getattr(self, "current_kernel_files", []), + }, + ) + # Add compilation failure to history so LLM learns from the error + if hasattr(self, "current_kernel_files") and self.current_kernel_files: + diff = self.compute_diff(self.current_kernel_files) + 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": 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", + { + "success": True, + "kernel_files": getattr(self, "current_kernel_files", []), + }, + ) return result def profile_pass(self) -> Result: @@ -103,7 +342,11 @@ def profile_pass(self) -> Result: # Log profiling results if hasattr(self, "_initial_profiler_results") and self._initial_profiler_results: self.get_logger().record( - "profile_pass_complete", {"profiler_results": self._initial_profiler_results, "top_n": self.top_n} + "profile_pass_complete", + { + "profiler_results": self._initial_profiler_results, + "top_n": self.top_n, + }, ) def instrument_pass(self) -> Result: @@ -117,7 +360,11 @@ def instrument_pass(self) -> Result: # Log instrumentation completion self.get_logger().record( - "instrument_pass_complete", {"success": True, "note": "Instrumentation pass completed via parent class"} + "instrument_pass_complete", + { + "success": True, + "note": "Instrumentation pass completed via parent class", + }, ) return Result( @@ -126,7 +373,12 @@ def instrument_pass(self) -> Result: error_report="The instrumentation is not implemented for memory access.", ) - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: + def optimize_pass( + self, + temperature: float = 0.0, + max_tokens: int = 3000, + target_kernel: str = None, + ) -> Result: """ Optimize the kernel to remove uncoalesced memory access via OpenAI API @@ -138,26 +390,14 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target Result: Optimized kernel as a file path """ super().optimize_pass() - llm_key = get_llm_api_key() system_prompt = ( - "You are a skilled GPU HIP programmer. Given a kernel," - " you will optimize it to remove uncoalesced memory access as much as possible" - " and provide a correct performant implementation. Do not modify" - " the kernel signature. Do not touch any other code, licenses, copyrights, or comments in the file." - " If you remove the copyright, your solution will be rejected." - " Do not include any markdown code blocks or text other than the code." + "You are a skilled GPU HIP programmer specializing in optimizing kernels " + "to improve memory coalescing and access patterns." ) - provider = self.provider - model = self.model - llm = LLM( - api_key=llm_key, - system_prompt=system_prompt, - model=model, - provider=provider, - logger=self.get_logger(), # Add logger here - ) + # Get LLM instance (initialized once per formula) + llm = self.get_llm(system_prompt) kernel = None kernel_file = None @@ -176,13 +416,19 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target ) if len(filtered_report_card) == 0: - return Result(success=False, error_report="No uncoalesced memory access found.") + logging.error("No uncoalesced memory access found.") + sys.exit(1) logging.debug(f"Filtered Report Card:\n{json.dumps(filtered_report_card, indent=4)}") kernel = filtered_report_card[0]["kernel"] self._parse_kernel_signature(kernel) + # Store baseline metrics on first run + if self.baseline_coalesced_pct is None: + self.baseline_coalesced_pct = filtered_report_card[0]["l1"]["coal"] + self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 + files = filtered_report_card[0]["source"]["files"] kernel_name = get_kernel_name(kernel) @@ -201,23 +447,15 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target kernel_file = file break if kernel_file is None: - return Result( - success=False, - error_report=f"Kernel file not found for kernel {kernel}.", - ) + logging.error(f"Kernel file not found for kernel {kernel}") + sys.exit(1) - user_prompt = ( - f"There is an uncoalesced memory access in the kernel {kernel} in the source code {unoptimized_file_content}." - f" Please fix the access pattern but do not change the semantics of the program." - " Do not remove any comments or licenses." - " Do not include any markdown code blocks or text other than the code." + # Build problem description + problem_description = ( + f"Uncoalesced memory access detected in kernel {kernel}. " + f"Please fix the access pattern to improve memory coalescing but do not change the semantics." ) - if self.current_summary is not None: - user_prompt += f"\n\nThe current summary is: {self.current_summary}" - cur_diff = self.compute_diff([kernel_file]) - user_prompt += f"\nThe diff between the current and initial code is: {cur_diff}" - self.previous_source_code = unoptimized_file_content if self.current_args: @@ -241,17 +479,65 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target return Result(success=False, error_report="Failed to extract the kernel file path.") logging.debug(f"System prompt: {system_prompt}") - logging.debug(f"LLM prompt: {user_prompt}") + logging.debug(f"Problem description: {problem_description}") self.current_kernel_files = [kernel_file] + try: - optimized_file_content = llm.ask(user_prompt).strip() + # Use DSPy signature with history - pass inputs as kwargs + # Build failure analysis from history + tracker_dict = self.optimization_tracker.to_dict() + failed_steps = [s for s in tracker_dict.get("steps", []) if not s.get("success", False)] + previous_failures_text = "" + if failed_steps: + previous_failures_text = "Previous failed attempts:\n" + for i, step in enumerate(failed_steps[-3:], 1): # Show last 3 failures + previous_failures_text += f"\nAttempt {i}: {step.get('report', 'Unknown error')}\n" + else: + previous_failures_text = "No previous failures yet. This is an early iteration." + + # Get latest diff if available + latest_diff_text = "" + if tracker_dict.get("steps"): + latest_step = tracker_dict["steps"][-1] + latest_diff_text = latest_step.get("diff", "No diff available") + else: + latest_diff_text = "No previous attempts yet." + + response = llm.ask( + signature=MemoryAccessOptimization, + answer_type=None, # Return full response object + record_meta="memory_access_optimization", + kernel_code=unoptimized_file_content, + code_preservation_rules="", # Field description contains the critical rules + output_format_rules="", # Field description contains the format requirements + problem_description=problem_description, + baseline_coalesced_pct=str(self.baseline_coalesced_pct), + baseline_time_ms=str(self.baseline_time_ms), + history=self.optimization_tracker.get_dspy_history(), + previous_failures=previous_failures_text, + latest_diff=latest_diff_text, + optimization_techniques="", # Field description contains the full content + amd_specific_optimizations="", # Field description contains the full content + optimization_categories="", # Field description contains the full content + ) + + # Extract optimized code from response + if hasattr(response, "optimized_code"): + optimized_file_content = response.optimized_code.strip() + else: + # Fallback for string response + optimized_file_content = str(response).strip() + optimized_file_content = self.postprocess_llm_code(optimized_file_content) # Log successful optimization self.get_logger().record( "optimization_success", - {"optimized_code_length": len(optimized_file_content), "kernel_file": kernel_file}, + { + "optimized_code_length": len(optimized_file_content), + "kernel_file": kernel_file, + }, ) with open(kernel_file, "w") as f: @@ -266,7 +552,10 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target ) except Exception as e: error_msg = f"An unexpected error occurred - {str(e)}" - self.get_logger().record("optimization_error", {"error": error_msg, "error_type": type(e).__name__}) + self.get_logger().record( + "optimization_error", + {"error": error_msg, "error_type": type(e).__name__}, + ) logging.error(error_msg) return Result(success=False, error_report=error_msg) @@ -290,8 +579,49 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) Result: Validation status """ result = super().correctness_validation_pass(self.current_kernel, self.current_args, accordo_absolute_tolerance) + + # Log correctness validation result if not result: self.current_summary = result.error_report + self.get_logger().record( + "correctness_validation_failed", + { + "success": False, + "error_report": result.error_report, + "kernel": self.current_kernel, + "args": self.current_args, + }, + ) + # Add correctness failure to history so LLM learns from the error + if hasattr(self, "current_kernel_files") and self.current_kernel_files: + diff = self.compute_diff(self.current_kernel_files) + with open(self.current_kernel_files[0], "r") as f: + failed_code = f.read() + + 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", + { + "success": True, + "kernel": self.current_kernel, + "args": self.current_args, + }, + ) return result def performance_validation_pass(self) -> Result: @@ -352,11 +682,7 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - if not success or speedup < 1: - self.current_summary = self.optimization_report - return Result(success=False, error_report=self.optimization_report) - - # Log performance validation results + # Log performance validation results (always, even if failed) self.get_logger().record( "performance_validation_complete", { @@ -373,16 +699,83 @@ def performance_validation_pass(self) -> Result: logging.info(self.optimization_report) - self.success = True - return Result(success=True, asset={"log": 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 + return Result(success=False, error_report=self.optimization_report) def write_results(self, output_file: str = None): """ - Writes the results to the output file. + Writes the results to the output file using the best optimization attempt. """ + # Restore best results for output + self._optimization_results = self.best_optimization_results + self.optimization_report = self.best_iteration_report + + for file in self.current_kernel_files: + with open(file, "w") as f: + f.write(self.best_kernel_code) + + # 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": 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.success}, + additional_results={ + "formula": "memoryAccess", + "success": self.success, + "optimization_history": self.optimization_tracker.to_dict(), + **metric_fields, + }, ) def summarize_previous_passes(self): diff --git a/src/intelliperf/formulas/swizzling.py b/src/intelliperf/formulas/swizzling.py index 68abc86c..93d6987a 100644 --- a/src/intelliperf/formulas/swizzling.py +++ b/src/intelliperf/formulas/swizzling.py @@ -77,6 +77,7 @@ def __init__( provider: str = "openai", in_place: bool = False, unittest_command: str = None, + num_attempts: int = 10, ): super().__init__( name, @@ -89,6 +90,7 @@ def __init__( provider, in_place, unittest_command, + num_attempts, ) # This temp option allows us to toggle if we want a full or partial instrumentation report @@ -116,7 +118,7 @@ def __init__( self.last_applied_diff = None self.initial_source_code = None - self.max_iterations = 10 + self.max_iterations = self.num_attempts self.best_l2_improvement = 0.0 # Start at baseline (no improvement) self.best_speedup = 1.0 # Start at 1.0x (no speedup) self.best_diff = "" @@ -126,6 +128,12 @@ def __init__( self.l2_improvement_history = [] self.gpu_spec = GPUSpec() + # Store metrics for structured reporting + self.best_unoptimized_l2_hr = None + self.best_optimized_l2_hr = None + self.best_unoptimized_time = None + self.best_optimized_time = None + # Removed local compute_diff; using Formula_Base.compute_diff instead def build_pass(self, validate_build_result=True) -> Result: @@ -167,7 +175,12 @@ def instrument_pass(self) -> Result: error_report="The instrumentation is not implemented for swizzling.", ) - def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target_kernel: str = None) -> Result: + def optimize_pass( + self, + temperature: float = 0.0, + max_tokens: int = 3000, + target_kernel: str = None, + ) -> Result: """ Optimize the kernel to improve l2 hit rate through block swizzling via two-stage LLM approach @@ -294,7 +307,10 @@ def optimize_pass(self, temperature: float = 0.0, max_tokens: int = 3000, target self.memory_analysis_done = True except Exception as e: logging.error(f"Failed to get memory analysis - {str(e)}") - return Result(success=False, error_report=f"Failed to get memory analysis - {str(e)}") + return Result( + success=False, + error_report=f"Failed to get memory analysis - {str(e)}", + ) history_prompt_part = "" if self.iteration_history: @@ -412,11 +428,17 @@ def correctness_validation_pass(self, accordo_absolute_tolerance: float = 1e-6) If no unittest_command is provided, skip validation (treat as success). """ if not self.unittest_command: - return Result(success=True, asset={"log": "No unittest_command provided; skipping correctness validation."}) + return Result( + success=True, + asset={"log": "No unittest_command provided; skipping correctness validation."}, + ) success, output = self._application.run_unit_test() if not success: - return Result(success=False, error_report=f"Unit test validation failed. Output:\n{output}") + return Result( + success=False, + error_report=f"Unit test validation failed. Output:\n{output}", + ) return Result(success=True, asset={"log": output}) def performance_validation_pass(self) -> Result: @@ -497,6 +519,11 @@ def performance_validation_pass(self) -> Result: self.best_diff = self.last_applied_diff self.best_iteration_report = self.optimization_report self.best_optimization_results = self._optimization_results + # Store metrics for structured reporting + self.best_unoptimized_l2_hr = unoptimized_l2_hit_rate + self.best_optimized_l2_hr = optimized_l2_hit_rate + self.best_unoptimized_time = unoptimized_time + self.best_optimized_time = optimized_time with open(self.current_kernel_files[0], "r") as f: self.best_kernel_code = f.read() # Mark as successful if we achieved any improvement @@ -520,9 +547,30 @@ def write_results(self, output_file: str = None): with open(file, "w") as f: f.write(self.best_kernel_code) + # Build structured metric fields + metric_fields = {} + if self.best_unoptimized_l2_hr is not None and self.best_optimized_l2_hr is not None: + metric_fields = { + "kernel_name": self.current_kernel, + "metric": "l2_hr_pct", # The counter we're optimizing (L2 cache hit rate) + "metric_name": "L2 Cache Hit Rate", # Human-readable name + "metric_before": self.best_unoptimized_l2_hr, + "metric_after": self.best_optimized_l2_hr, + "time_before_ms": ( + self.best_unoptimized_time / 1e6 if self.best_unoptimized_time else 0 + ), # Convert ns to ms + "time_after_ms": ( + self.best_optimized_time / 1e6 if self.best_optimized_time else 0 + ), # Convert ns to ms + } + super().write_results( output_file=output_file, - additional_results={"formula": "swizzling", "success": self.success}, + additional_results={ + "formula": "swizzling", + "success": self.success, + **metric_fields, + }, ) def summarize_previous_passes(self): From 9abf7b4007dae7e64e33a04565339be89fc65b1e Mon Sep 17 00:00:00 2001 From: Copilot <198982749+Copilot@users.noreply.github.com> Date: Wed, 15 Oct 2025 21:38:22 -0700 Subject: [PATCH 06/14] Implement Dynamic Versioning using setuptools-scm (#152) Co-authored-by: copilot-swe-agent[bot] <198982749+Copilot@users.noreply.github.com> Co-authored-by: mawad-amd <112003944+mawad-amd@users.noreply.github.com> Co-authored-by: Muhammad Awad --- pyproject.toml | 13 ++++++- src/intelliperf/__init__.py | 7 ++++ src/intelliperf/formulas/formula_base.py | 3 ++ tests/test_version.py | 47 ++++++++++++++++++++++++ 4 files changed, 68 insertions(+), 2 deletions(-) create mode 100644 tests/test_version.py diff --git a/pyproject.toml b/pyproject.toml index dc375f01..3f08b319 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -1,6 +1,9 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + [project] name = "intelliperf" -version = "0.0.1" +dynamic = ["version"] description = "Automated bottleneck detection and solution orchestration" authors = [ { name = "Muhammad Awad", email = "muhaawad@amd.com" }, @@ -25,7 +28,7 @@ include = ["intelliperf", "accordo"] intelliperf = "intelliperf.__main__:main" [build-system] -requires = ["setuptools>=61.0"] +requires = ["setuptools>=61", "wheel", "setuptools-scm>=8"] build-backend = "setuptools.build_meta" [tool.rocprofiler-compute] @@ -101,3 +104,9 @@ line-ending = "auto" [tool.ruff.lint.pydocstyle] convention = "google" # Use Google-style docstrings + +# ---- setuptools-scm versioning ---- +[tool.setuptools_scm] +version_scheme = "post-release" # .postN after last tag +local_scheme = "node-and-date" # add commit hash (e.g. +gabc1234) and date (e.g. +20250914) +fallback_version = "0.0.0" # used if git metadata unavailable diff --git a/src/intelliperf/__init__.py b/src/intelliperf/__init__.py index 3a1a8052..eeebbdf0 100644 --- a/src/intelliperf/__init__.py +++ b/src/intelliperf/__init__.py @@ -21,3 +21,10 @@ # OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE # SOFTWARE. ################################################################################ + +from importlib.metadata import PackageNotFoundError, version + +try: + __version__ = version("intelliperf") +except PackageNotFoundError: + __version__ = "unknown" diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 7fcc053c..35001319 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -40,6 +40,7 @@ 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 intelliperf import __version__ from intelliperf.core.application import Application from intelliperf.core.logger import Logger from intelliperf.utils.env import get_accordo_path @@ -643,11 +644,13 @@ def write_results( # create a new json contining optimized and unoptimized results if diagnose_only: results = { + "version": __version__, "initial": self._initial_profiler_results, **additional_results, } else: results = { + "version": __version__, "optimized": self._optimization_results, "initial": self._initial_profiler_results, "report_message": self.optimization_report, diff --git a/tests/test_version.py b/tests/test_version.py new file mode 100644 index 00000000..6b55d818 --- /dev/null +++ b/tests/test_version.py @@ -0,0 +1,47 @@ +"""Test dynamic versioning with setuptools-scm.""" + +import re + +import tomllib + + +def test_version_format(): + """Test that the package version follows the expected format.""" + try: + from importlib.metadata import version + + pkg_version = version("intelliperf") + # Version should match pattern: X.Y.Z or X.Y.Z.postN+hash.date + # Examples: "0.0.0", "0.0.0.post2+g07ded6f.d20251016", "1.2.3" + pattern = r"^\d+\.\d+\.\d+(\.post\d+(\+g[0-9a-f]+\.d\d{8})?)?$" + assert re.match(pattern, pkg_version), f"Version {pkg_version} doesn't match expected format" + except ImportError: + # Package not installed, skip test + pass + + +def test_fallback_version(): + """Test that fallback version is set correctly in config.""" + with open("pyproject.toml", "rb") as f: + config = tomllib.load(f) + + assert "setuptools_scm" in config["tool"], "setuptools_scm configuration missing" + assert config["tool"]["setuptools_scm"]["fallback_version"] == "0.0.0", "Fallback version should be 0.0.0" + + +def test_dynamic_version_config(): + """Test that version is configured as dynamic.""" + with open("pyproject.toml", "rb") as f: + config = tomllib.load(f) + + assert "version" in config["project"].get("dynamic", []), "Version should be in dynamic list" + assert "version" not in config["project"], "Static version should not be set" + + +def test_build_requirements(): + """Test that setuptools-scm is in build requirements.""" + with open("pyproject.toml", "rb") as f: + config = tomllib.load(f) + + build_reqs = config["build-system"]["requires"] + assert any("setuptools-scm" in req for req in build_reqs), "setuptools-scm should be in build requirements" From 95bd0697e9573a82da4ff456a5288c920bb1d121 Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Sun, 2 Nov 2025 23:00:29 -0800 Subject: [PATCH 07/14] Unify counters report (#155) Co-authored-by: github-actions[bot] --- .gitignore | 2 + pyproject.toml | 4 +- src/accordo/python/communicate.py | 140 ++++++++-- src/intelliperf/formulas/atomic_contention.py | 178 ++++++------- src/intelliperf/formulas/bank_conflict.py | 225 ++++++++-------- src/intelliperf/formulas/formula_base.py | 241 ++++++++++++++++-- src/intelliperf/formulas/memory_access.py | 217 +++++++--------- src/intelliperf/formulas/swizzling.py | 27 +- 8 files changed, 629 insertions(+), 405 deletions(-) diff --git a/.gitignore b/.gitignore index 41639bab..5132677b 100644 --- a/.gitignore +++ b/.gitignore @@ -14,3 +14,5 @@ external/ intelliperf_env/ trace/ .build/ + +.rocprofv3/ \ No newline at end of file diff --git a/pyproject.toml b/pyproject.toml index 3f08b319..99bc8197 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -12,9 +12,9 @@ authors = [ license = { text = "MIT" } readme = "README.md" -requires-python = ">=3.10" +requires-python = ">=3.9" -# Python dependencies +# Python dependencies dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==2.6.27", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]", "rpds-py"] [tool.setuptools] diff --git a/src/accordo/python/communicate.py b/src/accordo/python/communicate.py index ba433b99..d24760ec 100644 --- a/src/accordo/python/communicate.py +++ b/src/accordo/python/communicate.py @@ -23,10 +23,13 @@ ################################################################################ import ctypes +import errno import logging +import math import os import stat import sys +import threading import time import ml_dtypes @@ -36,6 +39,34 @@ from hip import memcpy_d2h, open_ipc_handle +def run_with_timeout(func, timeout_seconds, *args, **kwargs): + """Cross-platform timeout wrapper using threading. + + Runs func in a thread and raises TimeoutError if it doesn't complete in time. + """ + result = [None] + exception = [None] + + def target(): + try: + result[0] = func(*args, **kwargs) + except Exception as e: + exception[0] = e + + thread = threading.Thread(target=target, daemon=True) + thread.start() + thread.join(timeout=timeout_seconds) + + if thread.is_alive(): + # Thread is still running - timeout occurred + raise TimeoutError(f"Operation timed out after {timeout_seconds} seconds") + + if exception[0] is not None: + raise exception[0] + + return result[0] + + def read_ipc_handles(args, ipc_file_name): count = sum(1 for arg in args if "*" in arg and "const" not in arg) @@ -45,7 +76,6 @@ def read_ipc_handles(args, ipc_file_name): while len(handles) < count: if not os.path.exists(ipc_file_name): - logging.debug("Waiting for IPC file...") time.sleep(0.1) continue @@ -71,18 +101,18 @@ def read_ipc_handles(args, ipc_file_name): size_value = int.from_bytes(size_data, byteorder="little") sizes.append(size_value) - logging.debug("Final IPC Handle (hex):") - for i in range(0, len(handle_np), 16): - chunk = handle_np[i : i + 16] - logging.debug(" ".join(f"{b:02x}" for b in chunk)) - - logging.debug(f"Corresponding Pointer Size: {size_value} bytes") + # Verbose IPC handle debugging (only when new handle received) + if logging.getLogger().isEnabledFor(logging.DEBUG): + logging.debug("Final IPC Handle (hex):") + for i in range(0, len(handle_np), 16): + chunk = handle_np[i : i + 16] + logging.debug(" ".join(f"{b:02x}" for b in chunk)) + logging.debug(f"Corresponding Pointer Size: {size_value} bytes") if len(handles) < count: - logging.debug(f"Waiting for {count - len(handles)} more IPC handles...") + # Don't spam logs in hot loop - removed logging.debug here time.sleep(0.1) - # logging.debug(f"Successfully read {len(handles)} IPC handles and sizes.") return handles, sizes @@ -91,27 +121,83 @@ def send_response(pipe_name): fifo.write("done\n") -def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30): +def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30, process_pid=None, baseline_time_ms=None): logging.debug(f"pipe_name: {pipe_name}") logging.debug(f"get_kern_arg_data args: {args}") logging.debug(f"ipc_file_name: {ipc_file_name}") - if not os.path.exists(pipe_name): - os.mkfifo(pipe_name) - os.chmod(pipe_name, stat.S_IRWXU | stat.S_IRWXG | stat.S_IRWXO) - - start_time = time.time() - with open(pipe_name, "rb") as fifo: # noqa: F841 - while True: - if time.time() - start_time > ipc_timeout_seconds: - raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data") - - try: - ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) - break - except Exception as e: - if time.time() - start_time > ipc_timeout_seconds: - raise TimeoutError(f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data: {str(e)}") - time.sleep(0.1) + + # Calculate dynamic timeout based on baseline performance + if baseline_time_ms is not None and baseline_time_ms > 0: + # 2x baseline, rounded up to next second, minimum 3 seconds + ipc_timeout_seconds = max(3, math.ceil(baseline_time_ms / 1000.0 * 2.0)) + logging.debug(f"Using dynamic timeout: {ipc_timeout_seconds}s (2x baseline of {baseline_time_ms}ms)") + else: + logging.debug(f"Using default timeout: {ipc_timeout_seconds}s (no baseline available)") + + def _do_ipc_work(): + """Inner function that does the actual IPC work - wrapped with timeout""" + fifo = None + try: + if not os.path.exists(pipe_name): + os.mkfifo(pipe_name) + os.chmod(pipe_name, stat.S_IRWXU | stat.S_IRWXG | stat.S_IRWXO) + + # Try to open the pipe, checking if the process is still alive + while fifo is None: + # Check if the process is still alive (crash detection, not timeout) + if process_pid is not None: + try: + os.kill(process_pid, 0) # Signal 0 just checks if process exists + except OSError: + raise RuntimeError( + f"Accordo process (PID {process_pid}) crashed or terminated before opening pipe. Check for segfaults or GPU memory access errors." + ) + + try: + # Try non-blocking open + fd = os.open(pipe_name, os.O_RDONLY | os.O_NONBLOCK) + fifo = os.fdopen(fd, "rb") + break + except OSError as e: + if e.errno == errno.ENXIO: # ENXIO - no writer connected yet + time.sleep(0.1) + continue + else: + raise + + # Read IPC handles + while True: + # Check if the process is still alive (crash detection, not timeout) + if process_pid is not None: + try: + os.kill(process_pid, 0) + except OSError: + raise RuntimeError( + f"Accordo process (PID {process_pid}) crashed or terminated during execution. Check for segfaults or GPU memory access errors." + ) + + try: + ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) + break + except Exception: + # For non-timeout exceptions, retry with a short sleep + time.sleep(0.1) + + return ipc_handles, ptr_sizes + + finally: + if fifo: + fifo.close() + + # Run IPC work with timeout wrapper (cross-platform) + try: + ipc_handles, ptr_sizes = run_with_timeout(_do_ipc_work, ipc_timeout_seconds) + except TimeoutError: + # Enhance timeout message with context + timeout_msg = f"Timeout after {ipc_timeout_seconds} seconds during IPC communication" + if baseline_time_ms is not None: + timeout_msg += f" (baseline: {baseline_time_ms}ms, 2x timeout: {ipc_timeout_seconds}s). Code may be correct but too slow to be worth profiling." + raise TimeoutError(timeout_msg) type_map = { "double*": ctypes.c_double, diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index 0ec0b732..18f690a0 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -24,7 +24,6 @@ import json import logging -import os import dspy @@ -33,12 +32,17 @@ OptimizationTracker, Result, filter_json_field, - get_kernel_name, ) class AtomicContentionOptimization(dspy.Signature): - """Optimize GPU kernel code to reduce atomic contention and improve performance.""" + """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. + """ kernel_code = dspy.InputField(desc="The current kernel source code that needs atomic contention optimization.") @@ -255,15 +259,15 @@ def __init__( self.bottleneck_report = None self.current_summary = None self.previous_source_code = None - self.success = False # Initialize optimization tracker - # Atomic contention optimization maximizes latency reduction (minimize atomic_lat) - # Automatically calculates latency_improvement from unoptimized_lat / optimized_lat + # 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) self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="latency_improvement", - maximize=True, + maximize=False, # False = minimize raw metric (latency) before_metric="unoptimized_lat", after_metric="optimized_lat", ) @@ -272,13 +276,6 @@ 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 @@ -360,8 +357,8 @@ def build_pass(self, validate_build_result=True) -> Result: "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, + "unoptimized_lat": self.baseline_atomic_latency or 0, + "optimized_lat": self.baseline_atomic_latency or 0, }, success=False, request=f"Optimize atomic contention in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", @@ -397,7 +394,8 @@ def optimize_pass( system_prompt = ( "You are a skilled GPU HIP programmer specializing in optimizing kernels " - "to reduce atomic contention and improve performance." + "to reduce atomic contention and improve performance. " + "You have terminal access to explore the codebase and edit files directly." ) # Get LLM instance (initialized once per formula) @@ -434,21 +432,7 @@ def optimize_pass( self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - kernel_file = None - unoptimized_file_content = None - for file in files: - project_dir = os.path.abspath(self._application.get_project_directory()) - file_path = os.path.abspath(file) - isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir - if os.path.exists(file) and isfile_in_project: - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - return Result(success=False, error_report="Kernel file not found.") + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) # Build problem description problem_description = ( @@ -600,15 +584,18 @@ 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": 0.0, - "unoptimized_time": 0, - "optimized_time": 0, - "unoptimized_latency": self.baseline_atomic_latency or 0, - "optimized_latency": self.baseline_atomic_latency or 0, + "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) }, success=False, request=f"Optimize atomic contention in kernel {getattr(self, 'current_kernel_signature', 'unknown')}", @@ -656,28 +643,51 @@ def performance_validation_pass(self) -> Result: optimized_time = optimized_results[0]["durations"]["ns"] optimized_metric = optimized_results[0][field][subfield] - success = optimized_metric < unoptimized_metric - speedup = unoptimized_time / optimized_time - metric_improvement = unoptimized_metric / optimized_metric if optimized_metric != 0 else 1 + # 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, + ) + + # Query calculated values from tracker + speedup = step.metrics.get("speedup", 1.0) + latency_improvement = step.metrics.get("latency_improvement", 1.0) - # Calculate cycle latency improvement percentage + # Calculate percentage improvement for readability 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 metric_improvement > 1: + if latency_improvement > 1: self.optimization_report += ( f"Atomic Contention Reduction: Successfully reduced atomic contention by " - f"{metric_improvement:.2f}x. " + f"{latency_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 / metric_improvement:.2f}x. " + f"{1 / latency_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). " ) @@ -696,17 +706,20 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Log performance validation results (always, even if failed) + # Update the step's report with the built message + step.report = self.optimization_report + + # Log performance validation results (using tracker's calculated values) self.get_logger().record( "performance_validation_complete", { - "success": success, + "success": step.success, "unoptimized_time_ns": unoptimized_time, "optimized_time_ns": optimized_time, "unoptimized_metric": unoptimized_metric, "optimized_metric": optimized_metric, "speedup": speedup, - "metric_improvement": metric_improvement, + "metric_improvement": latency_improvement, "cycle_latency_improvement": cycle_latency_improvement, "optimization_report": self.optimization_report, }, @@ -714,42 +727,6 @@ 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 @@ -759,27 +736,34 @@ def write_results(self, output_file: str = None): """ Writes the results to the output file using the best optimization attempt. """ - # Restore best results for output - self._optimization_results = self.best_optimization_results - self.optimization_report = self.best_iteration_report - - for file in self.current_kernel_files: - with open(file, "w") as f: - f.write(self.best_kernel_code) + # 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) # Extract metrics from best optimization step - best_step = self.optimization_tracker.to_dict().get("best_step", {}) - metrics = best_step.get("metrics", {}) + metrics = self.optimization_tracker.get_best_metrics() + + # 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 # 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": 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 + "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 } # Include optimization history in results @@ -787,7 +771,7 @@ def write_results(self, output_file: str = None): output_file=output_file, additional_results={ "formula": "atomicContention", - "success": self.success, + "success": self.optimization_tracker.is_successful(), "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index b096d768..fed70ec7 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -35,7 +35,6 @@ OptimizationTracker, Result, filter_json_field, - get_kernel_name, ) from intelliperf.utils.process import capture_subprocess_output from intelliperf.utils.regex import generate_ecma_regex_from_list @@ -267,15 +266,15 @@ def __init__( self.bottleneck_report = None self.current_summary = None self.previous_source_code = None - self.success = False # Initialize optimization tracker - # Bank conflict optimization maximizes conflict reduction (minimize bank conflicts) - # Automatically calculates conflict_improvement from unoptimized_conflicts / optimized_conflicts + # 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) self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="conflict_improvement", - maximize=True, + maximize=False, # False = minimize raw metric (conflicts) before_metric="unoptimized_conflicts", after_metric="optimized_conflicts", ) @@ -284,13 +283,6 @@ 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. @@ -320,21 +312,22 @@ 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": 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, - ) + 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, + ) else: self.get_logger().record( "build_pass_success", @@ -522,22 +515,7 @@ def optimize_pass( self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - kernel_file = None - - unoptimized_file_content = None - for file in files: - project_dir = os.path.abspath(self._application.get_project_directory()) - file_path = os.path.abspath(file) - isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir - if os.path.exists(file) and isfile_in_project: - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - return Result(success=False, error_report="Kernel file not found.") + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) # Build problem description problem_description = ( @@ -686,21 +664,24 @@ 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}" - 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, - ) + 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, + ) else: self.get_logger().record( "correctness_validation_success", @@ -734,50 +715,79 @@ def performance_validation_pass(self) -> Result: optimized_time = optimized_results[0]["durations"]["ns"] optimized_conflicts = optimized_results[0]["lds"]["bc"] - success = optimized_conflicts < unoptimized_conflicts - speedup = unoptimized_time / optimized_time - conflict_improvement = unoptimized_conflicts / optimized_conflicts if optimized_conflicts != 0 else 1 + # 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 conflict_improvement_percentage = ( - (unoptimized_conflicts - optimized_conflicts) / unoptimized_conflicts if unoptimized_conflicts != 0 else 0 - ) * 100 + (unoptimized_conflicts - optimized_conflicts) / unoptimized_conflicts * 100 + if unoptimized_conflicts > 0 + else 0 + ) + # Build report using tracker's calculated values self.optimization_report = "" # Format the conflict improvement message - if conflict_improvement_percentage > 1: + if conflict_improvement > 1: self.optimization_report += ( f"Bank Conflict Reduction: Successfully reduced shared memory bank conflicts by " - f"{conflict_improvement_percentage:.1f}%. " + f"{conflict_improvement:.2f}x. " f"Conflict ratio improved from {unoptimized_conflicts:.1f} to {optimized_conflicts:.1f} " - f"(lower values indicate fewer conflicts and better performance). " + f"({conflict_improvement_percentage:.1f}% reduction - lower values mean fewer conflicts). " ) else: self.optimization_report += ( f"Bank Conflict Increase: Bank conflicts increased by " - f"{abs(conflict_improvement_percentage):.1f}%. " + f"{1 / conflict_improvement:.2f}x. " f"Conflict ratio worsened from {unoptimized_conflicts:.1f} to {optimized_conflicts:.1f} " - f"(higher values indicate more conflicts and worse performance). " + f"({abs(conflict_improvement_percentage):.1f}% increase - higher values mean more conflicts). " ) # 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 / 1_000_000:.2f}ms to {optimized_time / 1_000_000:.2f}ms " + f"reduced from {unoptimized_time / 1e6:.2f}ms to {optimized_time / 1e6:.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 / 1_000_000:.2f}ms to {optimized_time / 1_000_000:.2f}ms " + f"increased from {unoptimized_time / 1e6:.2f}ms to {optimized_time / 1e6:.2f}ms " f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Log performance validation results (always, even if failed) + # Update the step's report with the built message + step.report = self.optimization_report + + # Log performance validation results (using tracker's calculated values) self.get_logger().record( "performance_validation_complete", { - "success": success, + "success": step.success, "unoptimized_time_ns": unoptimized_time, "optimized_time_ns": optimized_time, "unoptimized_conflicts": unoptimized_conflicts, @@ -791,42 +801,6 @@ 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 @@ -836,27 +810,34 @@ def write_results(self, output_file: str = None): """ Writes the results to the output file using the best optimization attempt. """ - # Restore best results for output - self._optimization_results = self.best_optimization_results - self.optimization_report = self.best_iteration_report - - for file in self.current_kernel_files: - with open(file, "w") as f: - f.write(self.best_kernel_code) + # 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) # Extract metrics from best optimization step - best_step = self.optimization_tracker.to_dict().get("best_step", {}) - metrics = best_step.get("metrics", {}) + metrics = self.optimization_tracker.get_best_metrics() + + # 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 # 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": 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 + "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 } # Include optimization history in results @@ -864,7 +845,7 @@ def write_results(self, output_file: str = None): output_file=output_file, additional_results={ "formula": "bankConflict", - "success": self.success, + "success": self.optimization_tracker.is_successful(), "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 35001319..8af86207 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -26,6 +26,7 @@ import json import logging import os +import subprocess import sys import time from abc import abstractmethod @@ -57,6 +58,7 @@ 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""" @@ -88,30 +90,75 @@ def __init__( # History messages for DSPy (stored as list of dicts) self.history_messages = [] + def is_successful(self) -> bool: + """Check if optimization was successful (any improvement over baseline)""" + if not self.best_step or not self.best_step.success: + return False + improvement = self.best_step.get_metric(self.primary_metric, 1.0) + speedup = self.best_step.get_metric("speedup", 1.0) + return improvement > 1.0 or speedup > 1.0 + + def get_best_code(self) -> str: + """Get the optimized code from the best step""" + if self.best_step: + return self.best_step.optimized_code + return "" + + def get_best_report(self) -> str: + """Get the optimization report from the best step""" + if self.best_step: + return self.best_step.report + return "" + + def get_best_metrics(self) -> dict: + """Get the metrics from the best step""" + if self.best_step: + return self.best_step.metrics + return {} + def add_step( self, diff: str, report: str, metrics: dict, - success: bool, + success: bool = None, request: str = "", optimized_code: str = "", ) -> OptimizationStep: - """Add step and auto-update best based on primary metric""" + """Add step and auto-calculate improvement, speedup, and success""" + # Auto-calculate speedup if time metrics are available + unopt_time = metrics.get("unoptimized_time", 0) + opt_time = metrics.get("optimized_time", 0) + if unopt_time > 0 and opt_time > 0: + metrics["speedup"] = unopt_time / opt_time + # Auto-calculate improvement if before/after metrics are configured if self.before_metric and self.after_metric: before = metrics.get(self.before_metric, 0) after = metrics.get(self.after_metric, 0) if before != 0: - improvement = after / before if after != 0 else 1.0 + # If maximize=False: we want to minimize raw metric (conflicts, latency), so improvement = before / after + # If maximize=True: we want to maximize raw metric (coalescing %, hit rate), so improvement = after / before + if not self.maximize: + improvement = before / after if after != 0 else 1.0 + else: + improvement = after / before if after != 0 else 1.0 metrics[self.primary_metric] = improvement + # Auto-determine success if not explicitly provided + if success is None: + improvement = metrics.get(self.primary_metric, 1.0) + speedup = metrics.get("speedup", 1.0) + # Success = metric improved (> 1.0) AND runtime didn't regress (>= 1.0) + success = (improvement > 1.0) and (speedup >= 1.0) + step = OptimizationStep( iteration=self.current_iteration, diff=diff, report=report, metrics=metrics, success=success, + optimized_code=optimized_code, ) self.steps.append(step) self.current_iteration += 1 @@ -146,15 +193,24 @@ def add_step( self.history_messages.append(history_entry) - # Auto-update best - if self.best_step is None: - self.best_step = step - else: - new_val = step.get_metric(self.primary_metric) - cur_val = self.best_step.get_metric(self.primary_metric) - - if (self.maximize and new_val > cur_val) or (not self.maximize and new_val < cur_val): + # Auto-update best based on primary metric (only for successful steps) + if success: + if self.best_step is None or not self.best_step.success: + # First successful step, or replacing a failed step self.best_step = step + else: + new_val = step.get_metric(self.primary_metric) + cur_val = self.best_step.get_metric(self.primary_metric) + + # With proper improvement calculation, we always want higher improvement values + if new_val > cur_val: + self.best_step = step + elif new_val == cur_val: + # Tie-breaker: prefer higher speedup (better runtime performance) + new_speedup = step.get_metric("speedup") + cur_speedup = self.best_step.get_metric("speedup") + if new_speedup > cur_speedup: + self.best_step = step return step @@ -331,6 +387,70 @@ 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. @@ -386,11 +506,67 @@ def build(self, validate_build_result=True): if success: return Result(success=success, asset={"log": result}) else: + # Filter compiler log to remove noise and keep only errors + filtered_log = self._filter_compiler_errors(result) return Result( success=success, - error_report="The application contains compiler errors. Here is the compiler log: " + result, + error_report="The application contains compiler errors. Here is the compiler log:\n" + filtered_log, ) + @staticmethod + def _filter_compiler_errors(compiler_log: str, max_lines: int = 50) -> str: + """Filter compiler log to show only errors and a summary, not all warnings.""" + lines = compiler_log.split("\n") + + errors = [] + notes = [] + gmake_errors = [] + warning_count = 0 + + for line in lines: + if ": error:" in line: + errors.append(line) + elif ": note:" in line: + notes.append(line) + elif line.strip().startswith("gmake") and ("***" in line or "Error" in line): + gmake_errors.append(line) + elif ": warning:" in line: + warning_count += 1 + + # Build filtered output + filtered = [] + + if warning_count > 0: + filtered.append(f"[{warning_count} warnings omitted - only showing errors]\n") + + if errors: + filtered.append("=== COMPILATION ERRORS ===") + for error in errors[:max_lines]: # Limit errors too + filtered.append(error) + if len(errors) > max_lines: + filtered.append(f"... and {len(errors) - max_lines} more errors") + + if notes: + filtered.append("\n=== NOTES ===") + for note in notes[:max_lines]: # Limit notes too + filtered.append(note) + + if gmake_errors: + filtered.append("\n=== BUILD FAILED ===") + for gmake_error in gmake_errors[-5:]: # Last 5 gmake errors + filtered.append(gmake_error) + + if not filtered: + # No errors found, maybe it's a different kind of failure + # Return first and last few lines + filtered.append("=== BUILD OUTPUT (TRUNCATED) ===") + filtered.extend(lines[:10]) + if len(lines) > 20: + filtered.append(f"\n... {len(lines) - 20} lines omitted ...\n") + filtered.extend(lines[-10:]) + + return "\n".join(filtered) + # ---------------------------------------------------- # Required methods to be implemented by child classes # ---------------------------------------------------- @@ -441,6 +617,11 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole 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}") @@ -482,19 +663,43 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole logging.debug(f"kernel_args: {kernel_args}") logging.debug(f"ipc_file_name: {ipc_file_name}") - original_dir = os.getcwd() - os.chdir(project_directory) - os.posix_spawn(binary, binary_with_args, env) - os.chdir(original_dir) + # 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 + ) + process_pid = process.pid + logging.debug(f"Launched {label} process with PID: {process_pid}") + try: - results[label] = get_kern_arg_data(pipe_name, kernel_args, ipc_file_name) + results[label] = get_kern_arg_data( + pipe_name, kernel_args, ipc_file_name, process_pid=process_pid, baseline_time_ms=baseline_time_ms + ) except TimeoutError as e: logging.error(f"Timeout while getting kernel argument data for {label}: {str(e)}") + process.kill() # Kill the hung process + + # 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." + ) + return Result( success=False, - error_report=f"Timeout while getting kernel argument data for {label}: {str(e)}. The code may have crashed.", + 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() diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index 3d5c2000..fc38afb7 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -24,7 +24,6 @@ import json import logging -import os import sys import dspy @@ -34,7 +33,6 @@ OptimizationTracker, Result, filter_json_field, - get_kernel_name, ) @@ -252,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 maximizes coalescing improvement - # Automatically calculates coal_improvement from unoptimized_coal / optimized_coal + # 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) self.optimization_tracker = OptimizationTracker( max_iterations=self.num_attempts, primary_metric="coal_improvement", - maximize=True, + maximize=True, # True = maximize raw metric (coalescing %) before_metric="unoptimized_coal", after_metric="optimized_coal", ) @@ -269,13 +267,6 @@ 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. @@ -305,21 +296,22 @@ 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": 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, - ) + 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, + ) else: self.get_logger().record( "build_pass_success", @@ -430,25 +422,7 @@ def optimize_pass( self.baseline_time_ms = filtered_report_card[0]["durations"]["ns"] / 1e6 files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - - logging.debug(f"Kernel name: {kernel_name}") - kernel_file = None - unoptimized_file_content = None - for file in files: - project_dir = os.path.abspath(self._application.get_project_directory()) - file_path = os.path.abspath(file) - isfile_in_project = os.path.commonpath([project_dir, file_path]) == project_dir - - if os.path.exists(file) and isfile_in_project: - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - logging.error(f"Kernel file not found for kernel {kernel}") - sys.exit(1) + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) # Build problem description problem_description = ( @@ -598,21 +572,24 @@ 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}" - 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, - ) + 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, + ) else: self.get_logger().record( "correctness_validation_success", @@ -646,26 +623,54 @@ def performance_validation_pass(self) -> Result: optimized_time = optimized_results[0]["durations"]["ns"] optimized_coal = optimized_results[0]["l1"]["coal"] - success = optimized_coal > unoptimized_coal - speedup = unoptimized_time / optimized_time - coal_improvement = optimized_coal / unoptimized_coal if optimized_coal != 0 else 1 + # 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() + + 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"(higher percentages indicate more efficient memory access patterns). " + f"({coal_pct_improvement:.1f}% increase - higher percentages indicate more efficient memory access). " ) 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"(lower percentages indicate less efficient memory access patterns). " + f"({coal_pct_decrease:.1f}% decrease - lower percentages indicate less efficient memory access). " ) # Format the performance improvement message @@ -682,11 +687,14 @@ def performance_validation_pass(self) -> Result: f"({(1 / speedup - 1) * 100:.1f}% slower)." ) - # Log performance validation results (always, even if failed) + # Update the step's report with the built message + step.report = self.optimization_report + + # Log performance validation results (using tracker's calculated values) self.get_logger().record( "performance_validation_complete", { - "success": success, + "success": step.success, "unoptimized_time_ns": unoptimized_time, "optimized_time_ns": optimized_time, "unoptimized_coal": unoptimized_coal, @@ -699,42 +707,6 @@ 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 @@ -744,27 +716,34 @@ def write_results(self, output_file: str = None): """ Writes the results to the output file using the best optimization attempt. """ - # Restore best results for output - self._optimization_results = self.best_optimization_results - self.optimization_report = self.best_iteration_report - - for file in self.current_kernel_files: - with open(file, "w") as f: - f.write(self.best_kernel_code) + # 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) # Extract metrics from best optimization step - best_step = self.optimization_tracker.to_dict().get("best_step", {}) - metrics = best_step.get("metrics", {}) + metrics = self.optimization_tracker.get_best_metrics() + + # 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 # 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": 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 + "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 } # Include optimization history in results @@ -772,7 +751,7 @@ def write_results(self, output_file: str = None): output_file=output_file, additional_results={ "formula": "memoryAccess", - "success": self.success, + "success": self.optimization_tracker.is_successful(), "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, diff --git a/src/intelliperf/formulas/swizzling.py b/src/intelliperf/formulas/swizzling.py index 93d6987a..cf5f5430 100644 --- a/src/intelliperf/formulas/swizzling.py +++ b/src/intelliperf/formulas/swizzling.py @@ -27,7 +27,6 @@ import logging import os import stat -import sys import dspy @@ -37,7 +36,6 @@ Formula_Base, Result, filter_json_field, - get_kernel_name, ) from intelliperf.utils.env import get_llm_api_key @@ -256,22 +254,8 @@ def optimize_pass( kernel = filtered_report_card[0]["kernel"] files = filtered_report_card[0]["source"]["files"] - kernel_name = get_kernel_name(kernel) - - logging.debug(f"Kernel name: {kernel_name}") - kernel_file = None - for file in files: - if os.path.exists(file): - with open(file, "r") as f: - unoptimized_file_content = f.read() - if kernel_name in unoptimized_file_content: - kernel_file = file - break - if kernel_file is None: - logging.error(f"Kernel file not found for kernel {kernel}") - sys.exit(1) - else: - logging.debug(f"Kernel file found for kernel {kernel}: {kernel_file}") + kernel_file, unoptimized_file_content = self.find_kernel_file(files, kernel) + logging.debug(f"Kernel file found for kernel {kernel}: {kernel_file}") # Stage 1: Memory access pattern analysis (only run once) if not self.memory_analysis_done: @@ -287,7 +271,7 @@ def optimize_pass( self.bottleneck_report = ( f"L2 Cache Locality Detection: IntelliPerf identified suboptimal L2 cache hit rate " - f"in kernel `{kernel_name}`. Poor cache locality occurs when " + f"in kernel `{kernel}`. Poor cache locality occurs when " f"blocks accessing related memory are scheduled to different XCDs with separate L2 caches, " f"reducing overall cache effectiveness." ) @@ -532,7 +516,10 @@ def performance_validation_pass(self) -> Result: if self.current_iteration < self.max_iterations: self.current_summary = self.optimization_report # Always return success=False to continue iterating - return Result(success=False, error_report=self.best_iteration_report) + error_msg = ( + self.best_iteration_report if self.best_iteration_report else "Continuing optimization iterations..." + ) + return Result(success=False, error_report=error_msg) return Result(success=True, asset={"log": self.best_iteration_report}) From 3e43ebff235b1e317178d34450c20e051d8e8b7d Mon Sep 17 00:00:00 2001 From: Muhammad Awad <112003944+mawad-amd@users.noreply.github.com> Date: Fri, 7 Nov 2025 20:14:49 -0800 Subject: [PATCH 08/14] Use new accordo APIs (#157) Co-authored-by: github-actions[bot] --- .../uncoalesced/uncoalesced.hip | 6 +- .../b2b_matrix_transpose.hip | 22 +- .../matrix_transpose/matrix_transpose.hip | 22 +- examples/bank_conflict/reduce/reduce.hip | 24 +- .../bank_conflict/synthetic/synthetic.hip | 16 +- .../transpose_scale_add.hip | 24 +- .../transpose_scale_add_templated.hip | 24 +- examples/basic/vector_add/vector_add.hip | 44 +- examples/contention/histogram/histogram.hip | 28 +- examples/contention/reduction/reduction.hip | 10 + .../reduction_optimized.hip | 10 + .../simple_reduction/simple_reduction.hip | 30 +- src/accordo/__init__.py | 120 ++++- src/accordo/_internal/__init__.py | 3 + src/accordo/_internal/codegen.py | 53 ++ src/accordo/_internal/hip_interop.py | 109 +++++ src/accordo/_internal/ipc/__init__.py | 3 + src/accordo/_internal/ipc/communication.py | 184 +++++++ src/accordo/config.py | 109 +++++ src/accordo/exceptions.py | 38 ++ src/accordo/result.py | 103 ++++ src/accordo/snapshot.py | 60 +++ src/accordo/validator.py | 453 ++++++++++++++++++ src/intelliperf/__main__.py | 4 + src/intelliperf/core/logger.py | 26 + src/intelliperf/formulas/atomic_contention.py | 6 +- src/intelliperf/formulas/bank_conflict.py | 6 +- src/intelliperf/formulas/diagnose_only.py | 19 +- src/intelliperf/formulas/formula_base.py | 204 ++++---- src/intelliperf/formulas/memory_access.py | 6 +- 30 files changed, 1541 insertions(+), 225 deletions(-) create mode 100644 src/accordo/_internal/__init__.py create mode 100644 src/accordo/_internal/codegen.py create mode 100644 src/accordo/_internal/hip_interop.py create mode 100644 src/accordo/_internal/ipc/__init__.py create mode 100644 src/accordo/_internal/ipc/communication.py create mode 100644 src/accordo/config.py create mode 100644 src/accordo/exceptions.py create mode 100644 src/accordo/result.py create mode 100644 src/accordo/snapshot.py create mode 100644 src/accordo/validator.py diff --git a/examples/access_pattern/uncoalesced/uncoalesced.hip b/examples/access_pattern/uncoalesced/uncoalesced.hip index e7703ae1..82ef614a 100644 --- a/examples/access_pattern/uncoalesced/uncoalesced.hip +++ b/examples/access_pattern/uncoalesced/uncoalesced.hip @@ -53,7 +53,7 @@ __global__ void matrix_transpose(const T* __restrict__ in, } int main() { using T = __hip_bfloat16; - + const int width = 1024; const int height = 1024; const int size = width * height; @@ -97,7 +97,7 @@ int main() { std::cout << (correct ? "Transpose correct βœ…" : "Transpose incorrect ❌") << "\n"; - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); return 0; } \ No newline at end of file diff --git a/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip b/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip index 126b698c..2f70de30 100644 --- a/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip +++ b/examples/bank_conflict/b2b_matrix_transpose/b2b_matrix_transpose.hip @@ -25,6 +25,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 __global__ void matrixTransposeShared_0(float* out, @@ -106,9 +116,9 @@ void runTranspose(int width, int height) { float* d_in; float* d_out; - hipMalloc(&d_in, width * height * sizeof(float)); - hipMalloc(&d_out, width * height * sizeof(float)); - hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_in, width * height * sizeof(float))); + hip_try(hipMalloc(&d_out, width * height * sizeof(float))); + hip_try(hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice)); dim3 blockSize(TILE_DIM, TILE_DIM); dim3 gridSize((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM); @@ -120,10 +130,10 @@ void runTranspose(int width, int height) { if (status != hipSuccess) { std::terminate(); } - hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost)); - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); } int main() { diff --git a/examples/bank_conflict/matrix_transpose/matrix_transpose.hip b/examples/bank_conflict/matrix_transpose/matrix_transpose.hip index 636fbd08..4d3ebe61 100644 --- a/examples/bank_conflict/matrix_transpose/matrix_transpose.hip +++ b/examples/bank_conflict/matrix_transpose/matrix_transpose.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 __global__ void matrixTransposeShared(float* out, @@ -62,9 +72,9 @@ void runTranspose(int width, int height) { float* d_in; float* d_out; - hipMalloc(&d_in, width * height * sizeof(float)); - hipMalloc(&d_out, width * height * sizeof(float)); - hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_in, width * height * sizeof(float))); + hip_try(hipMalloc(&d_out, width * height * sizeof(float))); + hip_try(hipMemcpy(d_in, h_in.data(), width * height * sizeof(float), hipMemcpyHostToDevice)); dim3 blockSize(TILE_DIM, TILE_DIM); dim3 gridSize((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM); @@ -74,10 +84,10 @@ void runTranspose(int width, int height) { if(status != hipSuccess){ std::terminate(); } - hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out.data(), d_out, width * height * sizeof(float), hipMemcpyDeviceToHost)); - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); } int main(int argc, char* argv[]) { diff --git a/examples/bank_conflict/reduce/reduce.hip b/examples/bank_conflict/reduce/reduce.hip index 4cb89c61..483d89fa 100644 --- a/examples/bank_conflict/reduce/reduce.hip +++ b/examples/bank_conflict/reduce/reduce.hip @@ -25,6 +25,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define BLOCK_SIZE 256 __global__ void reduce_kernel(const float *d_in, float *d_out, int n) { @@ -61,22 +71,22 @@ int main() { } float *d_in = nullptr, *d_out = nullptr; - hipMalloc((void **)&d_in, size); - hipMalloc((void **)&d_out, sizeof(float) * (numElements / BLOCK_SIZE)); + hip_try(hipMalloc((void **)&d_in, size)); + hip_try(hipMalloc((void **)&d_out, sizeof(float) * (numElements / BLOCK_SIZE))); - hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice); + hip_try(hipMemcpy(d_in, h_in, size, hipMemcpyHostToDevice)); int gridSize = (numElements + BLOCK_SIZE - 1) / BLOCK_SIZE; reduce_kernel<<>>(d_in, d_out, numElements); - hipDeviceSynchronize(); + hip_try(hipDeviceSynchronize()); - hipMemcpy(h_out, d_out, sizeof(float) * gridSize, hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out, d_out, sizeof(float) * gridSize, hipMemcpyDeviceToHost)); printf("First block sum: %f\n", h_out[0]); // Free resources. - hipFree(d_in); - hipFree(d_out); + hip_try(hipFree(d_in)); + hip_try(hipFree(d_out)); free(h_in); free(h_out); diff --git a/examples/bank_conflict/synthetic/synthetic.hip b/examples/bank_conflict/synthetic/synthetic.hip index f3af0c91..0b1852b1 100644 --- a/examples/bank_conflict/synthetic/synthetic.hip +++ b/examples/bank_conflict/synthetic/synthetic.hip @@ -24,6 +24,16 @@ SOFTWARE. #include #include + +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + #define BLOCK_SIZE 256 #define LDS_SIZE 256 @@ -52,20 +62,20 @@ int main() { // Allocate memory on the device float* d_out; - hipMalloc(&d_out, BLOCK_SIZE * sizeof(float)); + hip_try(hipMalloc(&d_out, BLOCK_SIZE * sizeof(float))); dim3 blockSize(BLOCK_SIZE); dim3 gridSize(1); hipLaunchKernelGGL(bankConflictKernel, gridSize, blockSize, 0, 0, d_out); // Copy the result back to the host - hipMemcpy(h_out, d_out, BLOCK_SIZE * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_out, d_out, BLOCK_SIZE * sizeof(float), hipMemcpyDeviceToHost)); for (int i = 0; i < BLOCK_SIZE; ++i) { std::cout << "h_out[" << i << "] = " << h_out[i] << std::endl; } - hipFree(d_out); + hip_try(hipFree(d_out)); return 0; } diff --git a/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip b/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip index 77f83a90..3926b017 100644 --- a/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip +++ b/examples/bank_conflict/transpose_scale_add/transpose_scale_add.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 __global__ void matrixTransposeShared(float* out, @@ -75,9 +85,9 @@ void runKernels(int width, int height) { std::iota(h_data.begin(), h_data.end(), 0.0f); float *d_data, *d_out; - hipMalloc(&d_data, num_elements * sizeof(float)); - hipMalloc(&d_out, num_elements * sizeof(float)); - hipMemcpy(d_data, h_data.data(), num_elements * sizeof(float), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_data, num_elements * sizeof(float))); + hip_try(hipMalloc(&d_out, num_elements * sizeof(float))); + hip_try(hipMemcpy(d_data, h_data.data(), num_elements * sizeof(float), hipMemcpyHostToDevice)); // --- 1. Scale all elements by 2.0 int blockSize = 256; @@ -93,11 +103,11 @@ void runKernels(int width, int height) { matrixTransposeShared<<>>(d_out, d_data, width, height); // --- Copy result back - hipDeviceSynchronize(); - hipMemcpy(h_out.data(), d_out, num_elements * sizeof(float), hipMemcpyDeviceToHost); + hip_try(hipDeviceSynchronize()); + hip_try(hipMemcpy(h_out.data(), d_out, num_elements * sizeof(float), hipMemcpyDeviceToHost)); - hipFree(d_data); - hipFree(d_out); + hip_try(hipFree(d_data)); + hip_try(hipFree(d_out)); // Print a few values std::cout << "Result (first 10 elements):\n"; diff --git a/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip b/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip index a0b3aa45..09c6ae0a 100644 --- a/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip +++ b/examples/bank_conflict/transpose_scale_add_templated/transpose_scale_add_templated.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + #define TILE_DIM 16 template @@ -76,9 +86,9 @@ void runKernels(int width, int height) { std::iota(h_data.begin(), h_data.end(), static_cast(0)); T *d_data, *d_out; - hipMalloc(&d_data, num_elements * sizeof(T)); - hipMalloc(&d_out, num_elements * sizeof(T)); - hipMemcpy(d_data, h_data.data(), num_elements * sizeof(T), hipMemcpyHostToDevice); + hip_try(hipMalloc(&d_data, num_elements * sizeof(T))); + hip_try(hipMalloc(&d_out, num_elements * sizeof(T))); + hip_try(hipMemcpy(d_data, h_data.data(), num_elements * sizeof(T), hipMemcpyHostToDevice)); // --- 1. Scale all elements int blockSize = 256; @@ -93,11 +103,11 @@ void runKernels(int width, int height) { dim3 gridDim((width + TILE_DIM - 1) / TILE_DIM, (height + TILE_DIM - 1) / TILE_DIM); matrixTransposeShared<<>>(d_out, d_data, width, height); - hipDeviceSynchronize(); - hipMemcpy(h_out.data(), d_out, num_elements * sizeof(T), hipMemcpyDeviceToHost); + hip_try(hipDeviceSynchronize()); + hip_try(hipMemcpy(h_out.data(), d_out, num_elements * sizeof(T), hipMemcpyDeviceToHost)); - hipFree(d_data); - hipFree(d_out); + hip_try(hipFree(d_data)); + hip_try(hipFree(d_out)); std::cout << "Result (first 10 elements):\n"; for (int i = 0; i < 10; ++i) diff --git a/examples/basic/vector_add/vector_add.hip b/examples/basic/vector_add/vector_add.hip index ff8da547..a88189f5 100644 --- a/examples/basic/vector_add/vector_add.hip +++ b/examples/basic/vector_add/vector_add.hip @@ -25,6 +25,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + __global__ void vector_add(const float* a, const float* b, float* c, size_t n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) @@ -46,38 +56,38 @@ int main() { } float *d_a, *d_b, *d_c; - hipMalloc(&d_a, size); - hipMalloc(&d_b, size); - hipMalloc(&d_c, size); + hip_try(hipMalloc(&d_a, size)); + hip_try(hipMalloc(&d_b, size)); + hip_try(hipMalloc(&d_c, size)); - hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice); - hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice); + hip_try(hipMemcpy(d_a, h_a, size, hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_b, h_b, size, hipMemcpyHostToDevice)); const int threadsPerBlock = 256; const int blocks = (N + threadsPerBlock - 1) / threadsPerBlock; // Create HIP events hipEvent_t start, stop; - hipEventCreate(&start); - hipEventCreate(&stop); + hip_try(hipEventCreate(&start)); + hip_try(hipEventCreate(&stop)); // Record start - hipEventRecord(start, 0); + hip_try(hipEventRecord(start, 0)); // Launch kernel hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threadsPerBlock), 0, 0, d_a, d_b, d_c, N); // Record stop - hipEventRecord(stop, 0); - hipEventSynchronize(stop); + hip_try(hipEventRecord(stop, 0)); + hip_try(hipEventSynchronize(stop)); // Calculate elapsed time float milliseconds = 0; - hipEventElapsedTime(&milliseconds, start, stop); + hip_try(hipEventElapsedTime(&milliseconds, start, stop)); std::cout << "Kernel execution time: " << milliseconds << " ms" << std::endl; // Copy result back to host - hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost); + hip_try(hipMemcpy(h_c, d_c, size, hipMemcpyDeviceToHost)); // Print some results for (size_t i = 0; i < 5; ++i) @@ -87,11 +97,11 @@ int main() { delete[] h_a; delete[] h_b; delete[] h_c; - hipFree(d_a); - hipFree(d_b); - hipFree(d_c); - hipEventDestroy(start); - hipEventDestroy(stop); + hip_try(hipFree(d_a)); + hip_try(hipFree(d_b)); + hip_try(hipFree(d_c)); + hip_try(hipEventDestroy(start)); + hip_try(hipEventDestroy(stop)); return 0; } diff --git a/examples/contention/histogram/histogram.hip b/examples/contention/histogram/histogram.hip index 25bcc996..0ace48f4 100644 --- a/examples/contention/histogram/histogram.hip +++ b/examples/contention/histogram/histogram.hip @@ -5,6 +5,16 @@ #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + using index_type = std::size_t; using real_type = double; @@ -75,17 +85,17 @@ int main(int argc, char* argv[]) { real_type* d_input; real_type* d_output; - hipMalloc(&d_labels, n * sizeof(index_type)); - hipMalloc(&d_flip, n * sizeof(index_type)); - hipMalloc(&d_input, n * sizeof(real_type)); - hipMalloc(&d_output, k * sizeof(real_type)); + hip_try(hipMalloc(&d_labels, n * sizeof(index_type))); + hip_try(hipMalloc(&d_flip, n * sizeof(index_type))); + hip_try(hipMalloc(&d_input, n * sizeof(real_type))); + hip_try(hipMalloc(&d_output, k * sizeof(real_type))); - hipMemcpy(d_labels, h_labels.data(), n * sizeof(index_type), hipMemcpyHostToDevice); - hipMemcpy(d_flip, h_flip.data(), n * sizeof(index_type), hipMemcpyHostToDevice); - hipMemcpy(d_input, h_input.data(), n * sizeof(real_type), hipMemcpyHostToDevice); - hipMemcpy(d_output, h_output.data(), k * sizeof(real_type), hipMemcpyHostToDevice); + hip_try(hipMemcpy(d_labels, h_labels.data(), n * sizeof(index_type), hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_flip, h_flip.data(), n * sizeof(index_type), hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_input, h_input.data(), n * sizeof(real_type), hipMemcpyHostToDevice)); + hip_try(hipMemcpy(d_output, h_output.data(), k * sizeof(real_type), hipMemcpyHostToDevice)); histogram(n, k, d_labels, d_flip, d_input, d_output); - hipDeviceSynchronize(); + hip_try(hipDeviceSynchronize()); } \ No newline at end of file diff --git a/examples/contention/reduction/reduction.hip b/examples/contention/reduction/reduction.hip index c6d2948e..78886216 100644 --- a/examples/contention/reduction/reduction.hip +++ b/examples/contention/reduction/reduction.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + __global__ void reduction_kernel(const float* input, float* result, std::size_t count) { const auto thread_id = threadIdx.x + blockIdx.x * blockDim.x; if (thread_id < count) { diff --git a/examples/contention/reduction_optimized/reduction_optimized.hip b/examples/contention/reduction_optimized/reduction_optimized.hip index 4772bbdf..72db0f73 100644 --- a/examples/contention/reduction_optimized/reduction_optimized.hip +++ b/examples/contention/reduction_optimized/reduction_optimized.hip @@ -27,6 +27,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + __global__ void reduction_kernel(const float* input, float* result, std::size_t count) { extern __shared__ float shared_data[]; const auto thread_id = threadIdx.x + blockIdx.x * blockDim.x; diff --git a/examples/contention/simple_reduction/simple_reduction.hip b/examples/contention/simple_reduction/simple_reduction.hip index 1339f9a0..92f9535b 100644 --- a/examples/contention/simple_reduction/simple_reduction.hip +++ b/examples/contention/simple_reduction/simple_reduction.hip @@ -28,6 +28,16 @@ SOFTWARE. #include #include +#define hip_try(expr) \ + do { \ + hipError_t err = (expr); \ + if (err != hipSuccess) { \ + const char* msg = hipGetErrorString(err); \ + throw std::runtime_error(std::string("HIP error: ") + msg); \ + } \ + } while (0) + + using data_t = double; __global__ void reduction_kernel(const data_t* input, data_t* result, std::size_t count) { @@ -45,33 +55,33 @@ int main() { data_t* d_input = nullptr; data_t* d_result = nullptr; - hipMalloc(&d_input, count * sizeof(data_t)); - hipMalloc(&d_result, sizeof(data_t)); + hip_try(hipMalloc(&d_input, count * sizeof(data_t))); + hip_try(hipMalloc(&d_result, sizeof(data_t))); std::vector h_input(count, 1); - hipMemcpy(d_input, h_input.data(), count * sizeof(data_t), hipMemcpyHostToDevice); - hipMemset(d_result, 0, sizeof(data_t)); + hip_try(hipMemcpy(d_input, h_input.data(), count * sizeof(data_t), hipMemcpyHostToDevice)); + hip_try(hipMemset(d_result, 0, sizeof(data_t))); std::cout << "input: " << d_input << std::endl; std::cout << "result: " << d_result << std::endl; reduction_kernel<<>>(d_input, d_result, count); - hipDeviceSynchronize(); + hip_try(hipDeviceSynchronize()); data_t h_result = 0; - hipMemcpy(&h_result, d_result, sizeof(data_t), hipMemcpyDeviceToHost); + hip_try(hipMemcpy(&h_result, d_result, sizeof(data_t), hipMemcpyDeviceToHost)); if (h_result != count) { std::cout << "Kernel failed. Expected: " << count << ", Got: " << h_result << "\n"; - hipFree(d_input); - hipFree(d_result); + hip_try(hipFree(d_input)); + hip_try(hipFree(d_result)); return -1; } else { std::cout << "Success!"; } std::cout << std::endl; - hipFree(d_input); - hipFree(d_result); + hip_try(hipFree(d_input)); + hip_try(hipFree(d_result)); return 0; } diff --git a/src/accordo/__init__.py b/src/accordo/__init__.py index 946f447d..419a21a5 100644 --- a/src/accordo/__init__.py +++ b/src/accordo/__init__.py @@ -1,29 +1,105 @@ -################################################################################ -# MIT License +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +"""Accordo: Automated side-by-side correctness validation for GPU kernels. -# Copyright (c) 2025 Advanced Micro Devices, Inc. All Rights Reserved. +Public API: + - ValidationConfig: Configuration for kernel validation + - KernelArg: Structured kernel argument representation + - Snapshot: Captured kernel argument data from binary execution + - ValidationResult: Result of validation with detailed metrics + - ArrayMismatch: Information about array validation failures + - Accordo: Main validator class for kernel validation + - Exceptions: AccordoError, AccordoBuildError, AccordoTimeoutError, etc. -# Permission is hereby granted, free of charge, to any person obtaining a copy -# of this software and associated documentation files (the "Software"), to deal -# in the Software without restriction, including without limitation the rights -# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell -# copies of the Software, and to permit persons to whom the Software is -# furnished to do so, subject to the following conditions: +Quick Example (one-off validation): + >>> from accordo import Accordo + >>> config = Accordo.Config( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... Accordo.KernelArg(name="result", type="double*"), + ... Accordo.KernelArg(name="input", type="const double*"), + ... ], + ... tolerance=1e-6 + ... ) + >>> validator = Accordo(config) + >>> result = validator.validate( + ... reference_binary=["./app_ref"], + ... optimized_binary=["./app_opt"], + ... working_directory=".", + ... baseline_time_ms=10.0 + ... ) -# The above copyright notice and this permission notice shall be included in all -# copies or substantial portions of the Software. +Efficient Example (multiple optimizations vs same reference): + >>> # Capture reference once (returns Snapshot object) + >>> ref_snapshot = validator.capture_snapshot( + ... binary=["./app_ref"], + ... working_directory=".", + ... timeout_seconds=30 + ... ) + >>> print(ref_snapshot) # Snapshot(binary='./app_ref', arrays=3, execution_time_ms=12.50) + >>> + >>> # Compare multiple optimizations + >>> for opt_binary in optimized_binaries: + ... opt_snapshot = validator.capture_snapshot( + ... binary=opt_binary, + ... working_directory=".", + ... timeout_seconds=60 + ... ) + ... result = validator.compare_snapshots(ref_snapshot, opt_snapshot) +""" -# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR -# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE -# SOFTWARE. -################################################################################ +# Public API exports +from .config import KernelArg, ValidationConfig +from .exceptions import ( + AccordoBuildError, + AccordoError, + AccordoProcessError, + AccordoTimeoutError, + AccordoValidationError, +) +from .result import ArrayMismatch, ValidationResult +from .snapshot import Snapshot +from .validator import Accordo as _Accordo -"""Accordo package for validation and verification.""" +# Version +__version__ = "0.2.0" -from .python import code_gen, communicate, hip, utils -__all__ = ["communicate", "code_gen", "utils", "hip"] +# Nest all classes under Accordo namespace +class Accordo(_Accordo): + """Main Accordo validator with nested classes for clean API. + + All Accordo components are accessible as Accordo.ClassName: + - Accordo.Config (ValidationConfig) + - Accordo.KernelArg + - Accordo.Snapshot + - Accordo.Result (ValidationResult) + - Accordo.ArrayMismatch + - Accordo.Error (AccordoError) + - Accordo.BuildError (AccordoBuildError) + - Accordo.TimeoutError (AccordoTimeoutError) + - Accordo.ProcessError (AccordoProcessError) + - Accordo.ValidationError (AccordoValidationError) + """ + + # Configuration + Config = ValidationConfig + KernelArg = KernelArg + + # Data structures + Snapshot = Snapshot + Result = ValidationResult + ArrayMismatch = ArrayMismatch + + # Exceptions + Error = AccordoError + BuildError = AccordoBuildError + TimeoutError = AccordoTimeoutError + ProcessError = AccordoProcessError + ValidationError = AccordoValidationError + + +# Public API +__all__ = [ + "Accordo", +] diff --git a/src/accordo/_internal/__init__.py b/src/accordo/_internal/__init__.py new file mode 100644 index 00000000..cfc5bea3 --- /dev/null +++ b/src/accordo/_internal/__init__.py @@ -0,0 +1,3 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +"""Internal implementation details for Accordo. Not part of public API.""" diff --git a/src/accordo/_internal/codegen.py b/src/accordo/_internal/codegen.py new file mode 100644 index 00000000..39f95f10 --- /dev/null +++ b/src/accordo/_internal/codegen.py @@ -0,0 +1,53 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""Code generation for Accordo C++ header files.""" + +import logging + + +def generate_kernel_header(args: list[str], additional_includes: list[str] = None) -> str: + """Generate C++ header file for kernel arguments. + + Args: + args: List of argument type strings (e.g., ["double*", "const float*", "int"]) + additional_includes: Optional list of additional include directives + + Returns: + Path to the generated header file + """ + if additional_includes is None: + additional_includes = [] + + header_path = "/tmp/KernelArguments.hpp" + member_names = [f"arg{i}" for i in range(len(args))] + members = ";\n ".join(f"{arg} {name}" for arg, name in zip(args, member_names)) + ";" + as_tuple_members = ", ".join(member_names) + + # Build includes section + includes_section = "#include \n" + includes_section += "#include // for float16\n" + includes_section += "#include // for bfloat16\n" + + if additional_includes: + includes_section += "\n// User-provided includes\n" + for include in additional_includes: + includes_section += f"#include {include}\n" + + header_content = f"""#pragma once +{includes_section} +struct KernelArguments {{ + {members} + + auto as_tuple() const {{ + return std::tie({as_tuple_members}); + }} +}}; +""" + + with open(header_path, "w") as header_file: + header_file.write(header_content) + + logging.debug(f"Generated header file: {header_path}") + logging.debug(f"Header content: {header_content}") + return header_path diff --git a/src/accordo/_internal/hip_interop.py b/src/accordo/_internal/hip_interop.py new file mode 100644 index 00000000..0dff7cc6 --- /dev/null +++ b/src/accordo/_internal/hip_interop.py @@ -0,0 +1,109 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""HIP interop functions for Accordo.""" + +import ctypes +import logging + +import numpy as np + +rt_path = "libamdhip64.so" +hip_runtime = ctypes.cdll.LoadLibrary(rt_path) + + +def hip_try(err): + """Check HIP error code and raise exception if error occurred.""" + if err != 0: + hip_runtime.hipGetErrorString.restype = ctypes.c_char_p + error_string = hip_runtime.hipGetErrorString(ctypes.c_int(err)).decode("utf-8") + raise RuntimeError(f"HIP error code {err}: {error_string}") + + +class hipIpcMemHandle_t(ctypes.Structure): + """HIP IPC memory handle structure.""" + + _fields_ = [("reserved", ctypes.c_char * 64)] + + +def open_ipc_handle(ipc_handle_data): + """Open a HIP IPC memory handle. + + Args: + ipc_handle_data: NumPy array of uint8 with 64 elements + + Returns: + Device pointer value + """ + ptr = ctypes.c_void_p() + hipIpcMemLazyEnablePeerAccess = ctypes.c_uint(1) + hip_runtime.hipIpcOpenMemHandle.argtypes = [ + ctypes.POINTER(ctypes.c_void_p), + hipIpcMemHandle_t, + ctypes.c_uint, + ] + + if isinstance(ipc_handle_data, np.ndarray): + if ipc_handle_data.dtype != np.uint8 or ipc_handle_data.size != 64: + logging.debug(f"ipc_handle_data.size: {ipc_handle_data.size}") + raise ValueError("ipc_handle_data must be a 64-element uint8 numpy array") + ipc_handle_bytes = ipc_handle_data.tobytes() + ipc_handle_data = (ctypes.c_char * 64).from_buffer_copy(ipc_handle_bytes) + else: + raise TypeError("ipc_handle_data must be a numpy.ndarray of dtype uint8 with 64 elements") + + raw_memory = ctypes.create_string_buffer(64) + ctypes.memset(raw_memory, 0x00, 64) + ipc_handle_struct = hipIpcMemHandle_t.from_buffer(raw_memory) + ipc_handle_data_bytes = bytes(ipc_handle_data) + ctypes.memmove(raw_memory, ipc_handle_data_bytes, 64) + + logging.debug("[ipc_handle_struct]:") + for i in range(0, len(ipc_handle_data_bytes), 16): + chunk = ipc_handle_data_bytes[i : i + 16] + logging.debug(" ".join(f"{b:02x}" for b in chunk)) + + hip_try( + hip_runtime.hipIpcOpenMemHandle( + ctypes.byref(ptr), + ipc_handle_struct, + hipIpcMemLazyEnablePeerAccess, + ) + ) + + return ptr.value + + +def memcpy_d2h(ptr, num_elements_to_copy, dtype): + """Copy data from device to host. + + Args: + ptr: Device pointer value + num_elements_to_copy: Number of elements to copy + dtype: C type of elements + + Returns: + NumPy array with copied data + """ + host_array = np.zeros(num_elements_to_copy, dtype=np.dtype(dtype)) + + hip_runtime.hipMemcpy.argtypes = [ + ctypes.c_void_p, + ctypes.c_void_p, + ctypes.c_size_t, + ctypes.c_int, + ] + bytes_to_copy = num_elements_to_copy * ctypes.sizeof(dtype) + logging.debug( + f"Copying {num_elements_to_copy * ctypes.sizeof(dtype)} bytes from {hex(ptr)} to {hex(host_array.ctypes.data)}" + ) + + hip_try( + hip_runtime.hipMemcpy( + ctypes.c_void_p(host_array.ctypes.data), + ctypes.c_void_p(ptr), + ctypes.c_size_t(bytes_to_copy), + 2, # hipMemcpyDeviceToHost + ) + ) + return host_array diff --git a/src/accordo/_internal/ipc/__init__.py b/src/accordo/_internal/ipc/__init__.py new file mode 100644 index 00000000..c0c05d33 --- /dev/null +++ b/src/accordo/_internal/ipc/__init__.py @@ -0,0 +1,3 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. +"""IPC communication modules for Accordo.""" diff --git a/src/accordo/_internal/ipc/communication.py b/src/accordo/_internal/ipc/communication.py new file mode 100644 index 00000000..c5f61265 --- /dev/null +++ b/src/accordo/_internal/ipc/communication.py @@ -0,0 +1,184 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""IPC communication for Accordo.""" + +import ctypes +import logging +import os +import stat +import time + +import ml_dtypes +import numpy as np + +from ..hip_interop import memcpy_d2h, open_ipc_handle + + +def read_ipc_handles(args, ipc_file_name): + """Read IPC handles and sizes from the IPC file. + + Args: + args: List of argument type strings + ipc_file_name: Path to the IPC file + + Returns: + Tuple of (handles, sizes) + """ + count = sum(1 for arg in args if "*" in arg and "const" not in arg) + + handles = [] + sizes = [] + handles_set = set() + + while len(handles) < count: + if not os.path.exists(ipc_file_name): + logging.debug("Waiting for IPC file...") + time.sleep(0.1) + continue + + with open(ipc_file_name, "rb") as file: + data = file.read() + + messages = data.split(b"BEGIN\n") + for message in messages: + if b"END\n" in message: + content = message.split(b"END\n")[0] + + if len(content) == 72: + handle_data = content[:64] + size_data = content[64:72] + + handle_np = np.frombuffer(handle_data, dtype=np.uint8) + handle_tuple = tuple(handle_np) + + if handle_tuple not in handles_set: + handles.append(handle_np) + handles_set.add(handle_tuple) + + size_value = int.from_bytes(size_data, byteorder="little") + sizes.append(size_value) + + logging.debug("Final IPC Handle (hex):") + for i in range(0, len(handle_np), 16): + chunk = handle_np[i : i + 16] + logging.debug(" ".join(f"{b:02x}" for b in chunk)) + + logging.debug(f"Corresponding Pointer Size: {size_value} bytes") + + if len(handles) < count: + logging.debug(f"Waiting for {count - len(handles)} more IPC handles...") + time.sleep(0.1) + + return handles, sizes + + +def send_response(pipe_name): + """Send completion response through named pipe.""" + with open(pipe_name, "w") as fifo: + fifo.write("done\n") + + +def get_kern_arg_data(pipe_name, args, ipc_file_name, ipc_timeout_seconds=30, process_pid=None, baseline_time_ms=None): + """Get kernel argument data via IPC. + + Args: + pipe_name: Path to the named pipe + args: List of argument type strings + ipc_file_name: Path to the IPC file + ipc_timeout_seconds: Timeout for IPC operations + process_pid: Process ID (for error messages) + baseline_time_ms: Baseline execution time (for dynamic timeout) + + Returns: + List of NumPy arrays with argument data + + Raises: + TimeoutError: If IPC operation times out + TypeError: If unsupported type encountered + """ + # Calculate dynamic timeout if baseline provided + if baseline_time_ms is not None: + # Use 2x baseline or minimum 3 seconds + dynamic_timeout = max(3.0, (baseline_time_ms / 1000.0) * 2.0) + ipc_timeout_seconds = dynamic_timeout + logging.debug(f"Using dynamic timeout: {ipc_timeout_seconds}s (2x baseline of {baseline_time_ms}ms)") + + logging.debug(f"pipe_name: {pipe_name}") + logging.debug(f"get_kern_arg_data args: {args}") + logging.debug(f"ipc_file_name: {ipc_file_name}") + + if not os.path.exists(pipe_name): + os.mkfifo(pipe_name) + os.chmod(pipe_name, stat.S_IRWXU | stat.S_IRWXG | stat.S_IRWXO) + + start_time = time.time() + with open(pipe_name, "rb") as fifo: # noqa: F841 + while True: + # Check if the process is still alive (crash detection, not timeout) + if process_pid is not None: + try: + os.kill(process_pid, 0) # Signal 0 just checks if process exists + except OSError: + raise RuntimeError( + f"Accordo process (PID {process_pid}) crashed or terminated during execution. " + "Check for segfaults or GPU memory access errors." + ) + + if time.time() - start_time > ipc_timeout_seconds: + timeout_msg = f"Timeout after {ipc_timeout_seconds} seconds during IPC communication" + if baseline_time_ms is not None: + timeout_msg += f" (baseline: {baseline_time_ms}ms, 2x timeout: {ipc_timeout_seconds}s). Code may be correct but too slow to be worth profiling." + raise TimeoutError(timeout_msg) + + try: + ipc_handles, ptr_sizes = read_ipc_handles(args, ipc_file_name) + break + except Exception as e: + if time.time() - start_time > ipc_timeout_seconds: + timeout_msg = f"Timeout after {ipc_timeout_seconds} seconds waiting for IPC data: {str(e)}" + if baseline_time_ms is not None: + timeout_msg += f" (baseline: {baseline_time_ms}ms, 2x timeout: {ipc_timeout_seconds}s). Code may be correct but too slow to be worth profiling." + raise TimeoutError(timeout_msg) + time.sleep(0.1) + + type_map = { + "double*": ctypes.c_double, + "float*": ctypes.c_float, + "int*": ctypes.c_int, + "std::size_t*": ctypes.c_size_t, + "__half*": np.float16, + "__hip_bfloat16*": ml_dtypes.bfloat16, + } + + results = [] + pointer_args = list(filter(lambda arg: "*" in arg and "const" not in arg, args)) + logging.debug(f"pointer_args: {pointer_args}") + + for handle, arg, array_size in zip(ipc_handles, pointer_args, ptr_sizes): + ptr = open_ipc_handle(handle) + logging.debug(f"Opened IPC Ptr: {ptr} (0x{ptr:x})") + arg_type = arg.split()[0] + logging.debug(f"arg_type: {arg_type}") + + if arg_type in type_map: + dtype = type_map[arg_type] + logging.debug(f"dtype: {dtype}") + + # Special handling for FP16 and bfloat16 + if arg_type == "__half*": + temp_array = memcpy_d2h(ptr, array_size // 2, ctypes.c_uint16) + host_array = np.frombuffer(temp_array, dtype=np.float16) + elif arg_type == "__hip_bfloat16*": + temp_array = memcpy_d2h(ptr, array_size // 2, ctypes.c_uint16) + host_array = np.frombuffer(temp_array, dtype=ml_dtypes.bfloat16) + else: + num_elements = array_size // ctypes.sizeof(dtype) + host_array = memcpy_d2h(ptr, num_elements, dtype) + else: + raise TypeError(f"Unsupported pointer type: {arg_type}") + + logging.debug(f"Received data from IPC ({arg_type}/{len(host_array)}): {host_array}") + results.append(host_array) + + return results diff --git a/src/accordo/config.py b/src/accordo/config.py new file mode 100644 index 00000000..b7c6d295 --- /dev/null +++ b/src/accordo/config.py @@ -0,0 +1,109 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""Configuration classes for Accordo validation.""" + +from dataclasses import dataclass, field +from typing import Union + + +@dataclass +class KernelArg: + """Represents a kernel argument with semantic information. + + Args: + name: Argument name (e.g., "result", "input", "count") + type: C/C++ type string (e.g., "double*", "const float*", "int") + + Examples: + >>> KernelArg(name="result", type="double*") + >>> KernelArg(name="input", type="const double*") + >>> KernelArg(name="count", type="unsigned long") + + Note: + Output arguments are identified by checking for "*" without "const" in the type. + This matches the existing IPC logic. + """ + + name: str + type: str + + @classmethod + def from_string(cls, type_str: str, name: str = None) -> "KernelArg": + """Create KernelArg from a plain type string (backward compatibility). + + Args: + type_str: C/C++ type string + name: Optional argument name (auto-generated if not provided) + + Returns: + KernelArg instance + """ + if name is None: + name = f"arg_{id(type_str)}" # Generate unique name + return cls(name=name, type=type_str) + + @classmethod + def from_dict(cls, d: dict) -> "KernelArg": + """Create KernelArg from a dictionary.""" + return cls(**d) + + +@dataclass +class ValidationConfig: + """Configuration for Accordo kernel validation. + + Args: + kernel_name: Name of the kernel to validate + kernel_args: List of kernel arguments (KernelArg, str, or dict) + additional_includes: C++ include directives for custom types + tolerance: Absolute tolerance for array comparison + timeout_multiplier: Timeout = baseline_time_ms * timeout_multiplier + log_level: Logging level ("DEBUG", "INFO", "WARNING", "ERROR") + + Examples: + >>> config = ValidationConfig( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... KernelArg(name="result", type="double*", direction="out"), + ... KernelArg(name="input", type="const double*", direction="in"), + ... "int" # Plain string (backward compat) + ... ], + ... additional_includes=['"my_types.h"', ''], + ... tolerance=1e-6 + ... ) + """ + + kernel_name: str + kernel_args: list[Union[KernelArg, str, dict]] + additional_includes: list[str] = field(default_factory=list) + tolerance: float = 1e-6 + timeout_multiplier: float = 2.0 + log_level: str = "WARNING" + + def __post_init__(self): + """Normalize kernel_args to KernelArg instances.""" + normalized_args = [] + for i, arg in enumerate(self.kernel_args): + if isinstance(arg, KernelArg): + normalized_args.append(arg) + elif isinstance(arg, str): + # Convert plain string to KernelArg + normalized_args.append(KernelArg.from_string(arg, name=f"arg{i}")) + elif isinstance(arg, dict): + # Convert dict to KernelArg + if "name" not in arg: + arg["name"] = f"arg{i}" + normalized_args.append(KernelArg.from_dict(arg)) + else: + raise TypeError(f"kernel_args must be KernelArg, str, or dict, got {type(arg)}") + + self.kernel_args = normalized_args + + def get_arg_types(self) -> list[str]: + """Get list of argument type strings (for backward compatibility).""" + return [arg.type for arg in self.kernel_args] + + def get_arg_names(self) -> list[str]: + """Get list of argument names.""" + return [arg.name for arg in self.kernel_args] diff --git a/src/accordo/exceptions.py b/src/accordo/exceptions.py new file mode 100644 index 00000000..5ead2700 --- /dev/null +++ b/src/accordo/exceptions.py @@ -0,0 +1,38 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""Custom exceptions for Accordo validation.""" + + +class AccordoError(Exception): + """Base exception for all Accordo errors.""" + + pass + + +class AccordoBuildError(AccordoError): + """Raised when Accordo C++ library fails to build.""" + + pass + + +class AccordoTimeoutError(AccordoError): + """Raised when a kernel execution exceeds the timeout.""" + + def __init__(self, message: str, timeout_seconds: float): + super().__init__(message) + self.timeout_seconds = timeout_seconds + + +class AccordoProcessError(AccordoError): + """Raised when the instrumented process crashes or fails.""" + + def __init__(self, message: str, exit_code: int = None): + super().__init__(message) + self.exit_code = exit_code + + +class AccordoValidationError(AccordoError): + """Raised when array validation fails.""" + + pass diff --git a/src/accordo/result.py b/src/accordo/result.py new file mode 100644 index 00000000..69bd649a --- /dev/null +++ b/src/accordo/result.py @@ -0,0 +1,103 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""Result classes for Accordo validation.""" + +from dataclasses import dataclass +from typing import Optional + +import numpy as np + + +@dataclass +class ArrayMismatch: + """Represents a mismatch between reference and optimized arrays. + + Args: + arg_index: Index of the argument that failed validation + arg_name: Name of the argument + arg_type: Type string of the argument + max_difference: Maximum absolute difference between arrays + mean_difference: Mean absolute difference between arrays + reference_sample: Sample values from reference array + optimized_sample: Sample values from optimized array + """ + + arg_index: int + arg_name: str + arg_type: str + max_difference: float + mean_difference: float + reference_sample: np.ndarray + optimized_sample: np.ndarray + + def __str__(self) -> str: + """Human-readable string representation.""" + return ( + f"Mismatch in arg '{self.arg_name}' ({self.arg_type}): " + f"max_diff={self.max_difference:.2e}, mean_diff={self.mean_difference:.2e}" + ) + + +@dataclass +class ValidationResult: + """Result of Accordo validation. + + Args: + is_valid: True if all arrays matched within tolerance + error_message: Error message if validation failed + mismatches: List of array mismatches + matched_arrays: Dictionary of successfully matched arrays + execution_time_ms: Execution times for reference and optimized kernels + timeout_used: Timeout value used (if applicable) + """ + + is_valid: bool + error_message: Optional[str] = None + mismatches: list[ArrayMismatch] = None + matched_arrays: dict[str, dict] = None + execution_time_ms: dict[str, float] = None + timeout_used: Optional[float] = None + + def __post_init__(self): + """Initialize default values.""" + if self.mismatches is None: + self.mismatches = [] + if self.matched_arrays is None: + self.matched_arrays = {} + if self.execution_time_ms is None: + self.execution_time_ms = {} + + @property + def num_arrays_validated(self) -> int: + """Total number of arrays validated (matched + mismatched).""" + return len(self.matched_arrays) + len(self.mismatches) + + @property + def num_mismatches(self) -> int: + """Number of array mismatches.""" + return len(self.mismatches) + + @property + def success_rate(self) -> float: + """Percentage of arrays that matched.""" + total = self.num_arrays_validated + if total == 0: + return 0.0 + return (len(self.matched_arrays) / total) * 100.0 + + def summary(self) -> str: + """Get a human-readable summary of validation results.""" + if self.is_valid: + return f"βœ“ Validation passed! {self.num_arrays_validated} arrays matched within tolerance." + else: + lines = [f"βœ— Validation failed: {self.error_message}"] + if self.mismatches: + lines.append(f"\nMismatched arrays ({len(self.mismatches)}):") + for mismatch in self.mismatches: + lines.append(f" - {mismatch}") + return "\n".join(lines) + + def __str__(self) -> str: + """String representation.""" + return self.summary() diff --git a/src/accordo/snapshot.py b/src/accordo/snapshot.py new file mode 100644 index 00000000..025f5020 --- /dev/null +++ b/src/accordo/snapshot.py @@ -0,0 +1,60 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""Snapshot: Represents captured kernel argument data from a binary execution.""" + +from dataclasses import dataclass +from typing import List + +import numpy as np + + +@dataclass +class Snapshot: + """Represents a captured snapshot of kernel argument data. + + Attributes: + arrays: List of numpy arrays containing kernel argument data + execution_time_ms: Time taken to execute and capture the snapshot (milliseconds) + binary: The binary command that was executed + working_directory: The directory where the binary was executed + + Example: + >>> snapshot = Snapshot( + ... arrays=[np.array([1, 2, 3]), np.array([4, 5, 6])], + ... execution_time_ms=12.5, + ... binary=["./my_app"], + ... working_directory="/path/to/project" + ... ) + >>> print(f"Captured {len(snapshot.arrays)} arrays in {snapshot.execution_time_ms}ms") + """ + + arrays: List[np.ndarray] + execution_time_ms: float + binary: List[str] + working_directory: str + + def __repr__(self) -> str: + """Pretty representation of snapshot.""" + binary_str = " ".join(self.binary) + return ( + f"Snapshot(binary='{binary_str}', " + f"arrays={len(self.arrays)}, " + f"execution_time_ms={self.execution_time_ms:.2f})" + ) + + def summary(self) -> str: + """Get a detailed summary of the snapshot.""" + binary_str = " ".join(self.binary) + lines = [ + "Snapshot Summary:", + f" Binary: {binary_str}", + f" Working Directory: {self.working_directory}", + f" Execution Time: {self.execution_time_ms:.2f}ms", + f" Number of Arrays: {len(self.arrays)}", + ] + + for i, arr in enumerate(self.arrays): + lines.append(f" Array {i}: shape={arr.shape}, dtype={arr.dtype}") + + return "\n".join(lines) diff --git a/src/accordo/validator.py b/src/accordo/validator.py new file mode 100644 index 00000000..3d68b00d --- /dev/null +++ b/src/accordo/validator.py @@ -0,0 +1,453 @@ +# SPDX-License-Identifier: MIT +# Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved. + +"""AccordoValidator: Main validation class for Accordo.""" + +import logging +import os +import signal +import subprocess +import time +from pathlib import Path +from typing import Optional + +import numpy as np + +from ._internal.codegen import generate_kernel_header +from ._internal.ipc.communication import get_kern_arg_data, send_response +from .config import ValidationConfig +from .exceptions import AccordoBuildError, AccordoProcessError, AccordoTimeoutError +from .result import ArrayMismatch, ValidationResult +from .snapshot import Snapshot + + +class _TimeoutException(Exception): + """Internal exception for timeout handling.""" + + pass + + +def _timeout_handler(signum, frame): + """Signal handler for timeout.""" + raise _TimeoutException("Operation timed out") + + +def _build_accordo(accordo_path: Path, parallel_jobs: int = 16) -> Path: + """Build Accordo C++ library. + + Args: + accordo_path: Path to Accordo directory + parallel_jobs: Number of parallel build jobs + + Returns: + Path to built library + + Raises: + AccordoBuildError: If build fails + """ + try: + # Configure with CMake + result = subprocess.run( + ["cmake", "-B", "build"], + cwd=accordo_path, + capture_output=True, + text=True, + check=True, + ) + logging.debug(f"CMake configure output: {result.stdout}") + + # Build + result = subprocess.run( + ["cmake", "--build", "build", "--parallel", str(parallel_jobs)], + cwd=accordo_path, + capture_output=True, + text=True, + check=True, + ) + logging.debug(f"CMake build output: {result.stdout}") + + lib_path = accordo_path / "build" / "lib" / "libaccordo.so" + if not lib_path.exists(): + raise AccordoBuildError(f"Library not found at {lib_path}") + + return lib_path + + except subprocess.CalledProcessError as e: + raise AccordoBuildError(f"Accordo build failed: {e.stderr}") + except Exception as e: + raise AccordoBuildError(f"Accordo build failed: {str(e)}") + + +def _validate_arrays(arr1: np.ndarray, arr2: np.ndarray, tolerance: float) -> bool: + """Validate two arrays are close within tolerance. + + Args: + arr1: First array + arr2: Second array + tolerance: Absolute tolerance + + Returns: + True if arrays match within tolerance + """ + return np.allclose(arr1, arr2, atol=tolerance, rtol=0) + + +class Accordo: + """Validator for GPU kernel correctness using Accordo. + + This class manages the entire validation pipeline: + - Building the Accordo C++ library + - Running instrumented processes + - Collecting kernel argument data via IPC + - Validating arrays match within tolerance + + Example: + >>> from accordo import AccordoValidator, ValidationConfig, KernelArg + >>> config = ValidationConfig( + ... kernel_name="my_kernel", + ... kernel_args=[ + ... KernelArg(name="result", type="double*", direction="out"), + ... KernelArg(name="input", type="const double*", direction="in"), + ... ], + ... tolerance=1e-6 + ... ) + >>> validator = AccordoValidator(config) + >>> result = validator.validate(reference_app, optimized_app) + >>> if result.is_valid: + ... print("Validation passed!") + """ + + def __init__( + self, + config: ValidationConfig, + accordo_path: Optional[Path] = None, + force_rebuild: bool = False, + parallel_jobs: int = 16, + ): + """Initialize AccordoValidator. + + Args: + config: Validation configuration + accordo_path: Path to Accordo directory (auto-detected if None) + force_rebuild: Force rebuild even if library exists + parallel_jobs: Number of parallel build jobs + """ + self.config = config + self.parallel_jobs = parallel_jobs + self._built = False + self._lib_path = None + + # Auto-detect accordo_path if not provided + if accordo_path is None: + # Try to find it relative to this file + accordo_dir = Path(__file__).parent + if (accordo_dir / "build").exists() or (accordo_dir / "CMakeLists.txt").exists(): + accordo_path = accordo_dir + else: + # Try environment variable + from intelliperf.utils.env import get_accordo_path + + accordo_path = Path(get_accordo_path()) + + self.accordo_path = Path(accordo_path) + logging.debug(f"Accordo path: {self.accordo_path}") + + # Build if forced or library doesn't exist + lib_path = self.accordo_path / "build" / "lib" / "libaccordo.so" + if force_rebuild or not lib_path.exists(): + logging.info("Building Accordo C++ library...") + self._lib_path = _build_accordo(self.accordo_path, parallel_jobs) + self._built = True + else: + self._lib_path = lib_path + self._built = True + + def capture_snapshot( + self, + binary: list[str], + working_directory: str = ".", + timeout_seconds: int = 30, + ) -> Snapshot: + """Capture a snapshot of kernel argument data from a binary execution. + + Args: + binary: Command to run binary (e.g., ["./app", "arg1"]) + working_directory: Directory to run binary from + timeout_seconds: Timeout for this capture + + Returns: + Snapshot object containing captured arrays and execution metadata + + Raises: + AccordoBuildError: If Accordo library not built + AccordoProcessError: If instrumented process crashes + AccordoTimeoutError: If execution exceeds timeout + + Note: + Binary must be pre-compiled. Accordo does not build applications. + """ + if not self._built: + raise AccordoBuildError("Accordo library not built") + + # Generate kernel header with additional includes + arg_types = self.config.get_arg_types() + generate_kernel_header(arg_types, self.config.additional_includes) + + # Wrap app run with timeout + old_handler = signal.signal(signal.SIGALRM, _timeout_handler) + signal.alarm(timeout_seconds) + + try: + start_time = time.time() + result_arrays = self._run_instrumented_app( + binary, working_directory, label="snapshot", baseline_time_ms=None + ) + signal.alarm(0) # Cancel alarm on success + execution_time_ms = (time.time() - start_time) * 1000 + + return Snapshot( + arrays=result_arrays, + execution_time_ms=execution_time_ms, + binary=binary, + working_directory=working_directory, + ) + except _TimeoutException: + signal.alarm(0) + logging.error(f"Snapshot capture timed out after {timeout_seconds}s") + raise AccordoTimeoutError( + f"Snapshot capture timed out after {timeout_seconds}s. This may indicate a GPU crash or hung process.", + timeout_seconds=timeout_seconds, + ) + except TimeoutError as e: + signal.alarm(0) + raise AccordoTimeoutError(f"Snapshot timeout: {str(e)}", timeout_seconds) + except RuntimeError as e: + signal.alarm(0) + raise AccordoProcessError(f"Process crashed during snapshot: {str(e)}") + finally: + signal.alarm(0) + signal.signal(signal.SIGALRM, old_handler) + + def compare_snapshots( + self, + reference_snapshot: Snapshot, + optimized_snapshot: Snapshot, + ) -> ValidationResult: + """Compare two snapshots and validate their arrays. + + Args: + reference_snapshot: Snapshot from reference binary (from capture_snapshot) + optimized_snapshot: Snapshot from optimized binary (from capture_snapshot) + + Returns: + ValidationResult with validation status and details + """ + results = { + "reference": reference_snapshot.arrays, + "optimized": optimized_snapshot.arrays, + } + execution_times = { + "reference": reference_snapshot.execution_time_ms, + "optimized": optimized_snapshot.execution_time_ms, + } + + return self._validate_results(results, execution_times) + + def validate( + self, + reference_binary: list[str], + optimized_binary: list[str], + working_directory: str = ".", + baseline_time_ms: Optional[float] = None, + ) -> ValidationResult: + """Validate optimized kernel against reference (convenience method). + + This is a convenience wrapper that captures both snapshots and compares them. + For better performance when validating multiple optimizations against the same + reference, use capture_snapshot() and compare_snapshots() directly. + + Args: + reference_binary: Command to run reference binary (e.g., ["./app", "arg1"]) + optimized_binary: Command to run optimized binary (e.g., ["./app_opt", "arg1"]) + working_directory: Directory to run binaries from + baseline_time_ms: Baseline execution time for dynamic timeout + + Returns: + ValidationResult with validation status and details + + Raises: + AccordoBuildError: If Accordo library not built + AccordoProcessError: If instrumented process crashes + AccordoTimeoutError: If execution exceeds timeout + + Note: + Both binaries must be pre-compiled. Accordo does not build applications. + """ + # Calculate timeouts + ref_timeout = 30 # Default for reference + if baseline_time_ms: + opt_timeout = int((baseline_time_ms * self.config.timeout_multiplier / 1000.0) + 30.0) + else: + opt_timeout = 30 + + # Capture snapshots + reference_snapshot = self.capture_snapshot(reference_binary, working_directory, ref_timeout) + optimized_snapshot = self.capture_snapshot(optimized_binary, working_directory, opt_timeout) + + # Compare + return self.compare_snapshots(reference_snapshot, optimized_snapshot) + + def _run_instrumented_app( + self, binary_cmd: list[str], working_directory: str, label: str, baseline_time_ms: Optional[float] = None + ) -> list[np.ndarray]: + """Run an instrumented application and collect kernel argument data. + + Args: + binary_cmd: Binary command with arguments (e.g., ["./app", "arg1"]) + working_directory: Directory to run the binary from + label: Label for this run ("reference" or "optimized") + baseline_time_ms: Baseline time for dynamic timeout + + Returns: + List of numpy arrays with kernel argument data + """ + timestamp = int(time.time() * 1000) # Use milliseconds for uniqueness + pipe_name = f"/tmp/kernel_pipe_{timestamp}_{label}" + ipc_file_name = f"/tmp/ipc_handle_{timestamp}_{label}.bin" + + # Clean up any existing files + for file_path in [pipe_name, ipc_file_name]: + if os.path.exists(file_path): + os.remove(file_path) + + # Set up environment + env = os.environ.copy() + env["HSA_TOOLS_LIB"] = str(self._lib_path) + env["KERNEL_TO_TRACE"] = self.config.kernel_name + + # Set log level + debug_level = logging.getLogger().getEffectiveLevel() + level_map = { + logging.WARNING: 0, + logging.INFO: 1, + logging.DEBUG: 2, + logging.NOTSET: 3, + } + env["ACCORDO_LOG_LEVEL"] = str(level_map.get(debug_level, 0)) + env["ACCORDO_PIPE_NAME"] = pipe_name + env["ACCORDO_IPC_OUTPUT_FILE"] = ipc_file_name + + # Launch process + logging.debug(f"Launching {label} process with PID for kernel {self.config.kernel_name}") + logging.debug(f"binary_cmd: {binary_cmd}") + logging.debug(f"working_directory: {working_directory}") + logging.debug(f"kernel_args: {self.config.get_arg_types()}") + logging.debug(f"ipc_file_name: {ipc_file_name}") + + original_dir = os.getcwd() + try: + os.chdir(working_directory) + process_pid = os.posix_spawn(binary_cmd[0], binary_cmd, env) + logging.debug(f"Launched {label} process with PID: {process_pid}") + finally: + os.chdir(original_dir) + + # Get kernel argument data via IPC + try: + result_arrays = get_kern_arg_data( + pipe_name, + self.config.get_arg_types(), + ipc_file_name, + process_pid=process_pid, + baseline_time_ms=baseline_time_ms, + ) + except TimeoutError: + # Kill the process if it timed out + try: + os.kill(process_pid, 9) + except (OSError, ProcessLookupError): + pass # Process already dead + raise + + # Send completion response + send_response(pipe_name) + + return result_arrays + + def _validate_results( + self, results: dict[str, list[np.ndarray]], execution_times: dict[str, float] + ) -> ValidationResult: + """Validate results from reference and optimized runs. + + Args: + results: Dictionary with "reference" and "optimized" array lists + execution_times: Execution times for each run + + Returns: + ValidationResult with validation status + """ + reference_arrays = results["reference"] + optimized_arrays = results["optimized"] + + if len(reference_arrays) != len(optimized_arrays): + return ValidationResult( + is_valid=False, + error_message=f"Array count mismatch: {len(reference_arrays)} vs {len(optimized_arrays)}", + execution_time_ms=execution_times, + ) + + mismatches = [] + matched_arrays = {} + + for i, (ref_arr, opt_arr) in enumerate(zip(reference_arrays, optimized_arrays)): + arg = self.config.kernel_args[i] + + if not _validate_arrays(ref_arr, opt_arr, self.config.tolerance): + # Array mismatch + diff = np.abs(ref_arr - opt_arr) + mismatch = ArrayMismatch( + arg_index=i, + arg_name=arg.name, + arg_type=arg.type, + max_difference=float(np.max(diff)), + mean_difference=float(np.mean(diff)), + reference_sample=ref_arr[:10] if len(ref_arr) > 10 else ref_arr, + optimized_sample=opt_arr[:10] if len(opt_arr) > 10 else opt_arr, + ) + mismatches.append(mismatch) + + logging.debug(f"Arrays at index {i} for arg '{arg.name}' ({arg.type}) are NOT close.") + logging.debug(f" Max difference: {mismatch.max_difference}") + logging.debug(f" Mean difference: {mismatch.mean_difference}") + else: + # Array matched + matched_arrays[arg.name] = { + "index": i, + "type": arg.type, + "size": len(ref_arr), + } + logging.debug(f"Arrays at index {i} for arg '{arg.name}' ({arg.type}) are close.") + + # Determine overall success + is_valid = len(mismatches) == 0 + + if is_valid: + return ValidationResult( + is_valid=True, + matched_arrays=matched_arrays, + execution_time_ms=execution_times, + ) + else: + # Build error message + error_lines = [f"Validation failed: {len(mismatches)} array(s) mismatched"] + for m in mismatches: + error_lines.append(f" - {m}") + error_message = "\n".join(error_lines) + + return ValidationResult( + is_valid=False, + error_message=error_message, + mismatches=mismatches, + matched_arrays=matched_arrays, + execution_time_ms=execution_times, + ) diff --git a/src/intelliperf/__main__.py b/src/intelliperf/__main__.py index bf2d0926..d9e5e9f0 100644 --- a/src/intelliperf/__main__.py +++ b/src/intelliperf/__main__.py @@ -258,6 +258,10 @@ def main(): optimizer = formula(**optimizer_args) + # Store trace path in logger for use in iteration logging + if hasattr(optimizer, "get_logger"): + optimizer.get_logger().trace_path = args.trace_path + # Helper function to flush logs if tracing is enabled def flush_logs_if_enabled(): if hasattr(optimizer, "get_logger") and args.trace_path: diff --git a/src/intelliperf/core/logger.py b/src/intelliperf/core/logger.py index d52f9170..65da6ed8 100644 --- a/src/intelliperf/core/logger.py +++ b/src/intelliperf/core/logger.py @@ -190,6 +190,32 @@ def get_run_summary(self) -> Dict[str, Any]: "events": self.buffer, } + def save_iteration_code(self, kernel_name: str, iteration_num: int, code_content: str) -> str: + """ + Save iteration code to a file in the trace directory. + + Args: + kernel_name: Name of the kernel being optimized + iteration_num: Iteration number + code_content: The code content to save + + Returns: + The path to the saved file + """ + # Get output directory from trace_path if available + if hasattr(self, "trace_path") and self.trace_path: + output_dir = os.path.abspath(self.trace_path) + os.makedirs(output_dir, exist_ok=True) + else: + output_dir = os.path.abspath(".") + + iteration_file = os.path.join(output_dir, f"{kernel_name}_iteration_{iteration_num}.hip") + with open(iteration_file, "w") as f: + f.write(code_content) + logging.info(f"Saved iteration {iteration_num} code to {iteration_file}") + + return iteration_file + def flush(self, output_file: Optional[str] = None) -> bool: """ Flush logs to output targets with error handling. diff --git a/src/intelliperf/formulas/atomic_contention.py b/src/intelliperf/formulas/atomic_contention.py index 18f690a0..d188ffc4 100644 --- a/src/intelliperf/formulas/atomic_contention.py +++ b/src/intelliperf/formulas/atomic_contention.py @@ -525,8 +525,8 @@ def optimize_pass( }, ) - with open(kernel_file, "w") as f: - f.write(optimized_file_content) + # Write and log optimized code immediately after LLM generation + self.write_and_log_optimized_code(kernel_file, optimized_file_content) logging.debug(f"Optimized file content: {optimized_file_content}") return Result( success=True, @@ -770,9 +770,9 @@ def write_results(self, output_file: str = None): super().write_results( output_file=output_file, additional_results={ + "optimization_history": self.optimization_tracker.to_dict(), "formula": "atomicContention", "success": self.optimization_tracker.is_successful(), - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) diff --git a/src/intelliperf/formulas/bank_conflict.py b/src/intelliperf/formulas/bank_conflict.py index fed70ec7..4c6eb9c8 100644 --- a/src/intelliperf/formulas/bank_conflict.py +++ b/src/intelliperf/formulas/bank_conflict.py @@ -607,8 +607,8 @@ def optimize_pass( }, ) - with open(kernel_file, "w") as f: - f.write(optimized_file_content) + # Write and log optimized code immediately after LLM generation + self.write_and_log_optimized_code(kernel_file, optimized_file_content) logging.debug(f"Optimized file content: {optimized_file_content}") return Result( success=True, @@ -844,9 +844,9 @@ def write_results(self, output_file: str = None): super().write_results( output_file=output_file, additional_results={ + "optimization_history": self.optimization_tracker.to_dict(), "formula": "bankConflict", "success": self.optimization_tracker.is_successful(), - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) diff --git a/src/intelliperf/formulas/diagnose_only.py b/src/intelliperf/formulas/diagnose_only.py index 43da42ce..0dc2db34 100644 --- a/src/intelliperf/formulas/diagnose_only.py +++ b/src/intelliperf/formulas/diagnose_only.py @@ -70,14 +70,25 @@ def optimize_pass(self, target_kernel: str = None): def compile_pass(self): return super().compile_pass() - def correctness_validation_pass(self): + def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tolerance: float = 1e-6): """ - Validate the optimized kernel by comparing the output with the reference kernel + Validate the optimized kernel by comparing the output with the reference kernel. + + Note: diagnose_only doesn't actually optimize, so this always returns success. + The method signature matches other formulas for API consistency. + + Args: + kernel: Kernel name (unused for diagnose_only) + kernel_args: Kernel arguments (unused for diagnose_only) + accordo_absolute_tolerance: Tolerance parameter (unused for diagnose_only) Returns: - Result: Validation status + Result: Validation status (always success for diagnose_only) """ - return super().correctness_validation_pass() + # diagnose_only doesn't optimize, so validation always succeeds + from intelliperf.formulas.formula_base import Result + + return Result(success=True, asset={"log": "diagnose_only: No optimization performed, validation skipped."}) def performance_validation_pass(self): return super().performance_validation_pass() diff --git a/src/intelliperf/formulas/formula_base.py b/src/intelliperf/formulas/formula_base.py index 8af86207..b592c9d1 100644 --- a/src/intelliperf/formulas/formula_base.py +++ b/src/intelliperf/formulas/formula_base.py @@ -26,7 +26,6 @@ import json import logging import os -import subprocess import sys import time from abc import abstractmethod @@ -38,13 +37,10 @@ import numpy as np import pandas as pd -from accordo.python.code_gen import generate_header -from accordo.python.communicate import get_kern_arg_data, send_response -from accordo.python.utils import run_subprocess +from accordo import Accordo from intelliperf import __version__ from intelliperf.core.application import Application from intelliperf.core.logger import Logger -from intelliperf.utils.env import get_accordo_path from intelliperf.utils.process import capture_subprocess_output, exit_on_fail @@ -349,6 +345,10 @@ def __init__( self._llm = None self._dspy_configured = False + # Accordo caching: validator and reference snapshot are created once and reused + self._accordo_validator = None + self._reference_snapshot = None + self.build() def get_logger(self) -> Logger: @@ -451,6 +451,32 @@ def find_kernel_file(self, files: list, kernel: str) -> tuple: return kernel_file, unoptimized_file_content + def write_and_log_optimized_code(self, kernel_file: str, optimized_code: str) -> None: + """ + Write optimized code to file and log it for future reference. + + This should be called immediately after LLM generates code, regardless of whether + it will compile or pass validation. + + Args: + kernel_file: Path to the kernel file to write + optimized_code: The optimized code content + """ + # Write the code to the kernel file + with open(kernel_file, "w") as f: + f.write(optimized_code) + + # Automatically detect iteration number from optimization tracker + iteration_num = len(self.optimization_tracker.steps) if hasattr(self, "optimization_tracker") else 0 + + # Log the iteration code immediately + kernel_name = get_kernel_name( + self.current_kernel_signature if hasattr(self, "current_kernel_signature") else "kernel" + ) + self.get_logger().save_iteration_code(kernel_name, iteration_num, optimized_code) + + logging.debug(f"Wrote and logged optimized code to {kernel_file} (iteration {iteration_num})") + def _parse_kernel_signature(self, kernel_signature: str): """ Parses a kernel signature to extract the kernel name and its arguments. @@ -596,7 +622,10 @@ def optimize_pass(self, target_kernel: str = None): @abstractmethod def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tolerance: float = 1e-6): """ - Validates the the application. + Validates the application using Accordo. + + Uses snapshot caching: reference app is captured once on first call, + then each optimized version is captured and compared to the cached reference. """ if self.unittest_command: success, output = self._application.run_unit_test() @@ -607,125 +636,70 @@ def correctness_validation_pass(self, kernel, kernel_args, accordo_absolute_tole ) return Result(success=True, asset={"log": output}) + # Build the optimized application first (Accordo doesn't build) self._application.build() - unoptimized_binary = self._application.get_app_cmd()[0] - optimized_binary = self._reference_app.get_app_cmd()[0] - - logging.debug(f"unoptimized_binary: {unoptimized_binary}") - logging.debug(f"optimized_binary: {optimized_binary}") - - accordo_directory = get_accordo_path() - - # Get baseline time if available (for dynamic timeout calculation) - baseline_time_ms = getattr(self, "baseline_time_ms", None) - if baseline_time_ms is not None: - logging.debug(f"Using baseline time for dynamic timeout: {baseline_time_ms}ms") - - results = {} - for app, label in zip([self._reference_app, self._application], ["unoptimized", "optimized"]): - logging.debug(f"Running accordo for {label}") - timestamp = int(time.time()) - pipe_name = f"/tmp/kernel_pipe_{timestamp}" - ipc_file_name = f"/tmp/ipc_handle_{timestamp}.bin" - - for file in [ipc_file_name, ipc_file_name]: - if os.path.exists(file): - os.remove(file) - generate_header(kernel_args) - - run_subprocess(["cmake", "-B", "build"], accordo_directory) - run_subprocess(["cmake", "--build", "build", "--parallel", "16"], accordo_directory) - lib = os.path.join(accordo_directory, "build", "lib", "libaccordo.so") - env = os.environ.copy() - env["HSA_TOOLS_LIB"] = lib - env["KERNEL_TO_TRACE"] = kernel - - # Get the debug level from logger and convert it - debug_level = logging.getLogger().getEffectiveLevel() - level_map = { - logging.WARNING: 0, # Warning - logging.INFO: 1, # Info - logging.DEBUG: 2, # Debug - logging.NOTSET: 3, # NOTEST - } - env["ACCORDO_LOG_LEVEL"] = str(level_map.get(debug_level, 0)) # Default to 0 (Warning) if level not found - env["ACCORDO_PIPE_NAME"] = pipe_name - env["ACCORDO_IPC_OUTPUT_FILE"] = ipc_file_name - - binary = app.get_app_cmd_without_args() - binary_with_args = app.get_app_cmd() - project_directory = app.get_project_directory() - logging.debug(f"binary: {binary}") - logging.debug(f"project_directory: {project_directory}") - logging.debug(f"kernel: {kernel}") - logging.debug(f"binary_with_args: {binary_with_args}") - logging.debug(f"kernel_args: {kernel_args}") - logging.debug(f"ipc_file_name: {ipc_file_name}") - - # Launch the process with Accordo and track its PID - process = subprocess.Popen( - binary_with_args, env=env, cwd=project_directory, stdout=subprocess.DEVNULL, stderr=subprocess.DEVNULL + # Create validator if not already created (and cache it) + if self._accordo_validator is None: + kernel_arg_objects = [ + Accordo.KernelArg(name=f"arg{i}", type=arg_type) for i, arg_type in enumerate(kernel_args) + ] + + config = Accordo.Config( + kernel_name=kernel, + kernel_args=kernel_arg_objects, + tolerance=accordo_absolute_tolerance, + timeout_multiplier=2.0, ) - process_pid = process.pid - logging.debug(f"Launched {label} process with PID: {process_pid}") + self._accordo_validator = Accordo(config) + logging.debug("Created and cached Accordo validator") + + # Capture reference snapshot if not already captured (and cache it) + if self._reference_snapshot is None: try: - results[label] = get_kern_arg_data( - pipe_name, kernel_args, ipc_file_name, process_pid=process_pid, baseline_time_ms=baseline_time_ms + reference_binary = self._reference_app.get_app_cmd() + working_dir = self._reference_app.get_project_directory() + + logging.debug("Capturing reference snapshot (will be cached)") + self._reference_snapshot = self._accordo_validator.capture_snapshot( + binary=reference_binary, working_directory=working_dir, timeout_seconds=30 ) - except TimeoutError as e: - logging.error(f"Timeout while getting kernel argument data for {label}: {str(e)}") - process.kill() # Kill the hung process + logging.debug(f"Reference snapshot captured in {self._reference_snapshot.execution_time_ms:.2f}ms") + except Exception as e: + logging.error(f"Failed to capture reference snapshot: {str(e)}") + return Result(success=False, error_report=f"Failed to capture reference snapshot: {str(e)}") + + # Capture optimized snapshot and compare with cached reference + try: + baseline_time = getattr(self, "baseline_time_ms", None) + optimized_binary = self._application.get_app_cmd() + working_dir = self._application.get_project_directory() + + # Calculate timeout for optimized + if baseline_time: + opt_timeout = int((baseline_time * 2.0 / 1000.0) + 30.0) + else: + opt_timeout = 30 - # Provide context-specific error message - if baseline_time_ms is not None: - error_msg = f"Optimization exceeded 2x baseline execution time for {label}: {str(e)}. Code may be correct but too slow to be worth profiling." - else: - error_msg = ( - f"Timeout while getting kernel argument data for {label}: {str(e)}. The code may have crashed." - ) + logging.debug("Capturing optimized snapshot") + optimized_snapshot = self._accordo_validator.capture_snapshot( + binary=optimized_binary, working_directory=working_dir, timeout_seconds=opt_timeout + ) + logging.debug(f"Optimized snapshot captured in {optimized_snapshot.execution_time_ms:.2f}ms") - return Result( - success=False, - error_report=error_msg, - ) - except RuntimeError as e: - logging.error(f"Accordo process crashed for {label}: {str(e)}") - return Result( - success=False, - error_report=f"Accordo process crashed for {label}: {str(e)}. This usually indicates a segfault or GPU memory access error in the kernel code.", - ) - send_response(pipe_name) - - # Wait for the process to finish - process.wait(timeout=5) - logging.debug(f"results unoptimized: {results['unoptimized']}") - logging.debug(f"results optimized: {results['optimized']}") - key0, key1 = results.keys() - for i in range(len(results[key0])): - if not validate_arrays(results[key0][i], results[key1][i], accordo_absolute_tolerance): - diff = np.abs(results[key0][i] - results[key1][i]) - logging.debug(f"Arrays at index {i} for '{key0}' and '{key1}' are NOT close.") - logging.debug(f" {key0}[{i}]: {results[key0][i]}") - logging.debug(f" {key1}[{i}]: {results[key1][i]}") - logging.debug(f" Difference: {diff}") - logging.debug(f" Max difference: {np.max(diff)}") + # Compare snapshots + validation_result = self._accordo_validator.compare_snapshots(self._reference_snapshot, optimized_snapshot) + if validation_result.is_valid: + logging.debug("Validation succeeded.") + return Result(success=True) else: - argument_name = kernel_args[i] - logging.debug( - f"Arrays at index {i} for '{key0}' and '{key1}' are close. The argument type is '{argument_name}'." - ) - for i in range(len(results[key0])): - if not validate_arrays(results[key0][i], results[key1][i], accordo_absolute_tolerance): - argument_name = kernel_args[i] - return Result( - success=False, - error_report=f"The optimized code output does not match the unoptimized code output. Values at index {i} for the '{argument_name}' pointer are NOT close.", - ) - logging.debug("Validation succeeded.") - return Result(success=True) + return Result(success=False, error_report=validation_result.error_message) + + except Exception as e: + logging.error(f"Accordo validation error: {str(e)}") + return Result(success=False, error_report=f"Accordo validation error: {str(e)}") @abstractmethod def performance_validation_pass(self): diff --git a/src/intelliperf/formulas/memory_access.py b/src/intelliperf/formulas/memory_access.py index fc38afb7..3822c5c2 100644 --- a/src/intelliperf/formulas/memory_access.py +++ b/src/intelliperf/formulas/memory_access.py @@ -514,8 +514,8 @@ def optimize_pass( }, ) - with open(kernel_file, "w") as f: - f.write(optimized_file_content) + # Write and log optimized code immediately after LLM generation + self.write_and_log_optimized_code(kernel_file, optimized_file_content) logging.debug(f"Optimized file content: {optimized_file_content}") return Result( success=True, @@ -750,9 +750,9 @@ def write_results(self, output_file: str = None): super().write_results( output_file=output_file, additional_results={ + "optimization_history": self.optimization_tracker.to_dict(), "formula": "memoryAccess", "success": self.optimization_tracker.is_successful(), - "optimization_history": self.optimization_tracker.to_dict(), **metric_fields, }, ) From 72d76c3ea6a0944539eaaeec87d73c5ecac4a62b Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 7 Nov 2025 22:38:28 -0600 Subject: [PATCH 09/14] Remove env for accordo --- src/accordo/validator.py | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/src/accordo/validator.py b/src/accordo/validator.py index 3d68b00d..bb77f085 100644 --- a/src/accordo/validator.py +++ b/src/accordo/validator.py @@ -139,15 +139,15 @@ def __init__( # Auto-detect accordo_path if not provided if accordo_path is None: - # Try to find it relative to this file + # Find it relative to this file (accordo package directory) 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()) + raise RuntimeError( + f"Could not find Accordo build directory. Expected at {accordo_dir}. " + "Please build Accordo first or specify accordo_path explicitly." + ) self.accordo_path = Path(accordo_path) logging.debug(f"Accordo path: {self.accordo_path}") From 83a32dfb7ead5563973655efd6d95ab91ca0c704 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 7 Nov 2025 22:38:41 -0600 Subject: [PATCH 10/14] Use new nexus --- src/intelliperf/core/application.py | 63 ++++++++++++++++++++--------- 1 file changed, 44 insertions(+), 19 deletions(-) diff --git a/src/intelliperf/core/application.py b/src/intelliperf/core/application.py index 7accae92..eeb87a64 100644 --- a/src/intelliperf/core/application.py +++ b/src/intelliperf/core/application.py @@ -34,7 +34,6 @@ from intelliperf.utils import process from intelliperf.utils.env import ( get_guided_tuning_path, - get_nexus_path, get_rocprofiler_path, ) from intelliperf.utils.process import capture_subprocess_output, exit_on_fail @@ -224,26 +223,52 @@ def clone(self): ) def collect_source_code(self): - nexus_directory = get_nexus_path() - lib = os.path.join(nexus_directory, "build", "lib", "libnexus.so") - env = os.environ.copy() - - with tempfile.TemporaryDirectory() as tmp: - json_result_file = os.path.join(tmp, "nexus_output.json") - - env["HSA_TOOLS_LIB"] = lib - env["NEXUS_LOG_LEVEL"] = "2" - env["NEXUS_OUTPUT_FILE"] = json_result_file - env["TRITON_ALWAYS_COMPILE"] = "1" - env["TRITON_DISABLE_LINE_INFO"] = "0" - capture_subprocess_output(self.get_app_cmd(), new_env=env, working_directory=self.get_project_directory()) - - if os.path.exists(json_result_file): - df_results = json.loads(open(json_result_file).read()) - else: - df_results = {"kernels": {}} + """ + Collect source code for GPU kernels using Nexus. + + Returns: + dict: Dictionary containing kernel information with assembly, HIP source, files, and line numbers + """ + try: + from nexus import Nexus + except ImportError: + logging.error("Nexus Python API not found. Please install it: pip install nexus") + return {"kernels": {}} + + try: + # Create Nexus tracer with warning log level + nexus = Nexus(log_level=2) + + # Additional environment for Triton kernels + triton_env = { + "TRITON_ALWAYS_COMPILE": "1", + "TRITON_DISABLE_LINE_INFO": "0", + } + + # Run the application and capture kernel trace + trace = nexus.run( + command=self.get_app_cmd(), + env=triton_env, + cwd=self.get_project_directory(), + ) + + # Convert trace to the expected format + df_results = {"kernels": {}} + for kernel in trace: + df_results["kernels"][kernel.name] = { + "assembly": kernel.assembly, + "hip": kernel.hip, + "files": kernel.files, + "lines": kernel.lines, + "signature": kernel.signature, + } + return df_results + except Exception as e: + logging.error(f"Failed to collect source code with Nexus: {e}") + return {"kernels": {}} + def get_binary_absolute_path(self): if self.get_project_directory() != "": binary = self.get_app_cmd_without_args() From 5b4fd44a9425e4d8c763aa8c0349c9ee8dd6b97d Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 7 Nov 2025 22:38:49 -0600 Subject: [PATCH 11/14] Remove unneded env functions --- src/intelliperf/utils/env.py | 8 -------- 1 file changed, 8 deletions(-) diff --git a/src/intelliperf/utils/env.py b/src/intelliperf/utils/env.py index 996d21ed..7a8a8a8f 100644 --- a/src/intelliperf/utils/env.py +++ b/src/intelliperf/utils/env.py @@ -34,18 +34,10 @@ def get_guided_tuning_path(): return (Path(__file__).resolve().parent / "../../../external/guided-tuning").resolve() -def get_accordo_path(): - return (Path(__file__).resolve().parent / "../../accordo").resolve() - - def get_rocprofiler_path(): return (Path(__file__).resolve().parent / "../../../external/rocprofiler-compute/src").resolve() -def get_nexus_path(): - return (Path(__file__).resolve().parent / "../../../external/nexus").resolve() - - def get_llm_api_key(): llm_key = os.environ.get("LLM_GATEWAY_KEY") if not llm_key: From 9faf2af7f9f683f391f6dd779393d8f896ddb9fb Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 7 Nov 2025 22:39:26 -0600 Subject: [PATCH 12/14] Update pytoml --- pyproject.toml | 27 +++++++++++++-------------- 1 file changed, 13 insertions(+), 14 deletions(-) diff --git a/pyproject.toml b/pyproject.toml index 99bc8197..b7339b98 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -15,7 +15,19 @@ readme = "README.md" requires-python = ">=3.9" # Python dependencies -dependencies = ["tomli", "tabulate", "ml_dtypes", "dspy==2.6.27", "pandas", "duckdb", "rich", "pytest", "litellm[proxy]", "rpds-py"] +dependencies = [ + "tomli", + "tabulate", + "ml_dtypes", + "dspy==2.6.27", + "pandas", + "duckdb", + "rich", + "pytest", + "litellm[proxy]", + "rpds-py", + "nexus @ git+https://github.com/AMDResearch/nexus.git@main", +] [tool.setuptools] package-dir = {"" = "src"} @@ -39,19 +51,6 @@ python3 -m pip install --ignore-installed blinker && python3 -m pip install -r requirements.txt """ - -[tool.nexus] -git = "https://github.com/AMDResearch/nexus.git" -branch = "main" -build_command = """ -export CC=${ROCM_PATH}/bin/hipcc -export CXX=${ROCM_PATH}/bin/hipcc -cmake -B build -DCMAKE_PREFIX_PATH=/opt/rocm\ - -DLLVM_INSTALL_DIR=/opt/rocm/llvm\ - -DCMAKE_BUILD_TYPE=Debug -cmake --build build --parallel 16 -""" - [project.optional-dependencies] dev = [ "ruff==0.3.0", From b7410da6cb615725a0236c0bb86c565c2c1745d9 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 7 Nov 2025 22:44:29 -0600 Subject: [PATCH 13/14] Inherit logging level from IntelliPerf to Nexus --- src/intelliperf/core/application.py | 17 +++++++++++++++-- 1 file changed, 15 insertions(+), 2 deletions(-) diff --git a/src/intelliperf/core/application.py b/src/intelliperf/core/application.py index eeb87a64..9876874e 100644 --- a/src/intelliperf/core/application.py +++ b/src/intelliperf/core/application.py @@ -236,8 +236,21 @@ def collect_source_code(self): return {"kernels": {}} try: - # Create Nexus tracer with warning log level - nexus = Nexus(log_level=2) + # Map Python logging level to Nexus log level + # Python: NOTSET=0, DEBUG=10, INFO=20, WARNING=30, ERROR=40, CRITICAL=50 + # Nexus: 0=none, 1=info, 2=warning, 3=error, 4=detail + current_level = logging.getLogger().getEffectiveLevel() + if current_level <= logging.DEBUG: + nexus_log_level = 4 # detail (most verbose) + elif current_level <= logging.INFO: + nexus_log_level = 1 # info + elif current_level <= logging.WARNING: + nexus_log_level = 2 # warning + else: + nexus_log_level = 0 # none + + # Create Nexus tracer with inherited log level + nexus = Nexus(log_level=nexus_log_level) # Additional environment for Triton kernels triton_env = { From 42294c7b41bf6d58a9d5b4977685ce37a8ad3a70 Mon Sep 17 00:00:00 2001 From: Muhammad Awad Date: Fri, 7 Nov 2025 22:48:16 -0600 Subject: [PATCH 14/14] Update Nexus installation command in error message --- src/intelliperf/core/application.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/intelliperf/core/application.py b/src/intelliperf/core/application.py index 9876874e..203d04d2 100644 --- a/src/intelliperf/core/application.py +++ b/src/intelliperf/core/application.py @@ -232,7 +232,7 @@ def collect_source_code(self): try: from nexus import Nexus except ImportError: - logging.error("Nexus Python API not found. Please install it: pip install nexus") + logging.error("Nexus Python API not found. Please install it: pip install git+https://github.com/AMDResearch/nexus.git@main") return {"kernels": {}} try: