From 293aa0ac5dc69d22726a7eea8cb3c09c1590550a Mon Sep 17 00:00:00 2001 From: Harald van Dijk Date: Fri, 29 Aug 2025 16:01:01 +0100 Subject: [PATCH] Switch to LLVM versions 20/21 as supported versions. LLVM 21.1.0 has been released, so by our policy of supporting the two most recent LLVM releases, we can now move from LLVM 19/20 to LLVM 20/21. --- .github/README.md | 12 +- .github/actions/do_build_llvm/action.yml | 4 +- .../actions/do_build_ock_artefact/action.yml | 4 +- .../do_build_run_sanitizers/action.yml | 2 +- .github/actions/run_opencl_cts/action.yml | 2 +- .github/actions/run_sycl_cts/action.yml | 2 +- .github/actions/setup_build/action.yml | 7 +- .github/dockerfiles/Dockerfile_22.04-aarch64 | 2 - .github/dockerfiles/Dockerfile_22.04-x86-64 | 7 +- .github/dockerfiles/Dockerfile_24.04-x86-64 | 5 +- .github/workflows/codeql.yml | 6 +- .github/workflows/create_llvm_artefacts.yml | 2 +- .../workflows/create_publish_artifacts.yml | 2 +- .github/workflows/planned_testing_caller.yml | 4 +- .../workflows/planned_testing_caller_19.yml | 71 -------- .../workflows/planned_testing_caller_20.yml | 4 +- .../workflows/planned_testing_caller_21.yml | 4 +- .../workflows/planned_testing_caller_main.yml | 4 +- .github/workflows/pr_tests_cache.yml | 4 +- .github/workflows/run_ock_external_tests.yml | 8 +- .github/workflows/run_ock_internal_tests.yml | 1 + .github/workflows/run_pr_tests_caller.yml | 4 +- CHANGELOG.md | 4 +- README.md | 2 +- cmake/ImportLLVM.cmake | 4 +- .../compiler/supported-llvm-versions.rst | 4 +- .../compiler/compiler_pipeline/CMakeLists.txt | 2 - .../utils/remove_address_spaces_pass.h | 43 ----- .../compiler_pipeline/source/builtin_info.cpp | 2 +- .../optimal_builtin_replacement_pass.cpp | 3 +- .../source/remove_address_spaces_pass.cpp | 152 ------------------ .../multi_llvm/include/multi_llvm/dibuilder.h | 101 ------------ .../include/multi_llvm/instructions.h | 3 - .../include/multi_llvm/instructions.inc | 14 +- .../multi_llvm/include/multi_llvm/intrinsic.h | 14 +- .../include/multi_llvm/loop_utils.h | 37 ----- .../include/multi_llvm/multi_llvm.h | 1 - .../source/base/include/base/target.h | 1 + .../source/base_module_pass_machinery.cpp | 1 - .../base/source/base_module_pass_registry.def | 3 - .../compiler/source/base/source/module.cpp | 2 - .../spirv-ll/include/spirv-ll/builder.h | 3 +- .../include/spirv-ll/builder_debug_info.h | 6 +- .../spirv-ll/source/builder_debug_info.cpp | 13 +- .../host/source/AddFloatingPointControl.cpp | 9 +- .../targets/host/source/HostPassMachinery.cpp | 6 - .../compiler/test/lit/passes/barriers-dbg.ll | 21 +-- modules/compiler/utils/source/lld_linker.cpp | 12 -- .../transform/packetization_helpers.cpp | 5 +- .../source/transform/pre_linearize_pass.cpp | 2 +- .../vecz/source/transform/scalarizer.cpp | 4 +- .../vecz/source/vector_target_info.cpp | 5 +- .../vecz/source/vector_target_info_arm.cpp | 5 +- .../vecz/source/vectorization_helpers.cpp | 4 +- .../packetize_mask_varying.ll | 8 +- .../vecz/test/lit/llvm/irreducible_loop.ll | 12 +- 56 files changed, 89 insertions(+), 580 deletions(-) delete mode 100644 .github/workflows/planned_testing_caller_19.yml delete mode 100644 modules/compiler/compiler_pipeline/include/compiler/utils/remove_address_spaces_pass.h delete mode 100644 modules/compiler/compiler_pipeline/source/remove_address_spaces_pass.cpp delete mode 100644 modules/compiler/multi_llvm/include/multi_llvm/dibuilder.h delete mode 100644 modules/compiler/multi_llvm/include/multi_llvm/loop_utils.h diff --git a/.github/README.md b/.github/README.md index 7e1157372e..cc31135ce0 100644 --- a/.github/README.md +++ b/.github/README.md @@ -20,9 +20,6 @@ create_llvm_artefacts.yml `Run planned testing`: planned_testing_caller.yml - description: runs planned_testing-style tests, called from an llvm version caller -`run planned tests for llvm 19`: planned_testing_caller_19.yml -- description: runs planned_tests for llvm 19 - `run planned tests for llvm 20`: planned_testing_caller_20.yml - description: runs planned_tests for llvm 20 @@ -60,8 +57,6 @@ create_llvm_artefacts.yml `CodeQL`: codeql.yml -`run planned tests for llvm 19`: planned_testing_caller_19.yml - `run planned tests for llvm 20`: planned_testing_caller_20.yml `run planned tests for llvm 21`: planned_testing_caller_21.yml @@ -74,8 +69,6 @@ create_llvm_artefacts.yml ### `workflow_dispatch:` workflows (manually runnable and available in forks) -`run planned tests for llvm 19`: planned_testing_caller_19.yml - `run planned tests for llvm 20`: planned_testing_caller_20.yml `run planned tests for llvm 21`: planned_testing_caller_21.yml @@ -114,7 +107,7 @@ The `publish docker images` workflow is configured to rebuild the containers whe ## LLVM artefact management -Planned_testing workflows each use a particular llvm artefact according to llvm version, OS and architecture (e.g. llvm 19/20/21, Ubuntu_24, x86_64). The specific version to use and branch to reference are contained in the .yml workflow definition. llvm artefacts can be installed, built or accessed as pre-built artefacts from Github cache. They are handled as follows: +Planned_testing workflows each use a particular llvm artefact according to llvm version, OS and architecture (e.g. llvm 20/21, Ubuntu_24, x86_64). The specific version to use and branch to reference are contained in the .yml workflow definition. llvm artefacts can be installed, built or accessed as pre-built artefacts from Github cache. They are handled as follows: - PR testing: llvm artefact is installed as needed. - Planned testing: a flag is set depending on whether the llvm version is stable. If the `use_llvm_github_cache` flag is set, cache is used with llvm artefact being built if required. If the flag is not set, the llvm artefact is always built. @@ -153,11 +146,10 @@ Planned_testing workflows are configured to run via `workflow_dispatch:` (manual ### Further tailoring of planned_testing workflows The following planned_testing workflows call `Run planned testing` (planned_testing_caller.yml) as a sub-workflow: ``` - run planned tests for llvm 19: planned_testing_caller_19.yml run planned tests for llvm 20: planned_testing_caller_20.yml run planned tests for llvm 21: planned_testing_caller_21.yml ``` -These workflows can be tailored to run specific llvm versions (e.g. 19), target lists (e.g. host_x86_64_linux) and test options (e.g. test_sanitizers), etc., by setting the scripted `inputs:` values to `Run planned testing` accordingly. See the planned_testing workflow .yml files for examples of current default values and tailoring options. With the exception of DPC++ PRs (see above) tailored values can only be set directly in workflow config and cannot currently be updated interactively on a per-run basis when called from the web interface (i.e. `workflow_dispatch:`). +These workflows can be tailored to run specific llvm versions (e.g. 20), target lists (e.g. host_x86_64_linux) and test options (e.g. test_sanitizers), etc., by setting the scripted `inputs:` values to `Run planned testing` accordingly. See the planned_testing workflow .yml files for examples of current default values and tailoring options. With the exception of DPC++ PRs (see above) tailored values can only be set directly in workflow config and cannot currently be updated interactively on a per-run basis when called from the web interface (i.e. `workflow_dispatch:`). ### Planned_testing workflows in forks Planned_testing workflows can also be run via `workflow_dispatch:` (manual event trigger) in forks. Examples can be found [in this fork](https://github.com/AERO-Project-EU/oneapi-construction-kit/actions?query=event%3Aworkflow_dispatch). diff --git a/.github/actions/do_build_llvm/action.yml b/.github/actions/do_build_llvm/action.yml index 9ecec3a520..8ae7054592 100644 --- a/.github/actions/do_build_llvm/action.yml +++ b/.github/actions/do_build_llvm/action.yml @@ -16,11 +16,11 @@ inputs: default: RelAssert llvm_version: type: string - description: llvm version - used for naming purposes e.g. 19,20,main" + description: llvm version - used for naming purposes e.g. 20,21,main" default: "" llvm_branch: type: string - description: "The actual llvm branch to check out e.g. release/19.x" + description: "The actual llvm branch to check out e.g. release/20.x" use_github_cache: type: boolean description: "Whether to use caching - will rebuild if not existing" diff --git a/.github/actions/do_build_ock_artefact/action.yml b/.github/actions/do_build_ock_artefact/action.yml index 182177d3c5..700e510f28 100644 --- a/.github/actions/do_build_ock_artefact/action.yml +++ b/.github/actions/do_build_ock_artefact/action.yml @@ -5,8 +5,8 @@ description: Action to build the oneapi-construction-kit as an artefact inputs: llvm_version: - description: 'llvm version we want to use (18-19)' - default: '19' + description: 'llvm version we want to use (20-21)' + default: '20' llvm_source: required: true description: 'source to get llvm from - one of install, cache or workflow id' diff --git a/.github/actions/do_build_run_sanitizers/action.yml b/.github/actions/do_build_run_sanitizers/action.yml index 021a05f271..b99141f218 100644 --- a/.github/actions/do_build_run_sanitizers/action.yml +++ b/.github/actions/do_build_run_sanitizers/action.yml @@ -11,7 +11,7 @@ inputs: type: string required: true llvm_branch: - description: 'The actual llvm branch to check out e.g. release/19.x' + description: 'The actual llvm branch to check out e.g. release/20.x' type: string required: true diff --git a/.github/actions/run_opencl_cts/action.yml b/.github/actions/run_opencl_cts/action.yml index d4ba363ece..b0f182cfb5 100644 --- a/.github/actions/run_opencl_cts/action.yml +++ b/.github/actions/run_opencl_cts/action.yml @@ -10,7 +10,7 @@ inputs: description: 'quick | full' default: "quick" llvm_version: - description: 'llvm major version (e.g 19,20, main) - to be used for llvm specific fails' + description: 'llvm major version (e.g 20,21,main) - to be used for llvm specific fails' required: true split_index: description: "optional split index so this run can be done with different splits, current only 0 or 1" diff --git a/.github/actions/run_sycl_cts/action.yml b/.github/actions/run_sycl_cts/action.yml index 99381ea345..22453e9752 100644 --- a/.github/actions/run_sycl_cts/action.yml +++ b/.github/actions/run_sycl_cts/action.yml @@ -8,7 +8,7 @@ inputs: description: '"opencl" or "native_cpu"' default: 'opencl' llvm_version: - description: 'llvm major version (e.g 19,20, main) - to be used for llvm specific fails' + description: 'llvm major version (e.g 20,21,main) - to be used for llvm specific fails' required: true split_index: description: "optional split index so this run can be done with different splits, current only 0 or 1" diff --git a/.github/actions/setup_build/action.yml b/.github/actions/setup_build/action.yml index 3d3e333279..c2d58997bd 100644 --- a/.github/actions/setup_build/action.yml +++ b/.github/actions/setup_build/action.yml @@ -9,7 +9,7 @@ inputs: description: 'llvm Build type (Release, RelAssert) - note we need to use RelAssert for the cache pattern matching' default: RelAssert llvm_version: - description: 'Major llvm version to use for fetching llvm cache e.g. 19' + description: 'Major llvm version to use for fetching llvm cache e.g. 20' ubuntu_version: description: 'Version of ubuntu used for cache retrieval and prerequisites' default: 22.04 @@ -90,8 +90,11 @@ runs: - name: install llvm native ubuntu # via "apt install" if: inputs.os == 'ubuntu' && inputs.llvm_source == 'install' shell: bash + env: + UBUNTU_NAME: "${{ inputs.ubuntu_version == '24.04' && 'noble' || 'jammy' }}" run: | set -x + echo "deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/$UBUNTU_NAME/ llvm-toolchain-$UBUNTU_NAME-${{ inputs.llvm_version }} main" | sudo tee /etc/apt/sources.list.d/llvm.list >/dev/null sudo apt update sudo apt install -y llvm-${{ inputs.llvm_version }}-dev liblld-${{ inputs.llvm_version }}-dev libclang-${{ inputs.llvm_version }}-dev @@ -104,7 +107,7 @@ runs: shell: pwsh env: # Map inputs.llvm_version to specific download release here. - LLVM_RELEASE: "${{ inputs.llvm_version == '20' && '20.1.2' || '19.1.7' }}" + LLVM_RELEASE: "${{ inputs.llvm_version == '21' && '21.1.0' || '20.1.8' }}" run: | Invoke-WebRequest https://github.com/llvm/llvm-project/releases/download/llvmorg-$env:LLVM_RELEASE/clang+llvm-$env:LLVM_RELEASE-x86_64-pc-windows-msvc.tar.xz -OutFile clang+llvm-$env:LLVM_RELEASE-x86_64-pc-windows-msvc.tar.xz 7z x clang+llvm-$env:LLVM_RELEASE-x86_64-pc-windows-msvc.tar.xz -so | 7z x -si -ttar diff --git a/.github/dockerfiles/Dockerfile_22.04-aarch64 b/.github/dockerfiles/Dockerfile_22.04-aarch64 index 874446de37..acb2cba559 100644 --- a/.github/dockerfiles/Dockerfile_22.04-aarch64 +++ b/.github/dockerfiles/Dockerfile_22.04-aarch64 @@ -10,8 +10,6 @@ RUN pip install cmakelint colorama lit RUN apt-get install --yes spirv-tools RUN apt-get install --yes wget gpg RUN wget -qO - https://apt.llvm.org/llvm-snapshot.gpg.key | gpg --dearmor - | tee /usr/share/keyrings/llvm-archive-keyring.gpg >/dev/null -RUN echo 'deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/jammy/ llvm-toolchain-jammy-19 main' | tee /etc/apt/sources.list.d/llvm.list >/dev/null -RUN echo 'deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/jammy/ llvm-toolchain-jammy-20 main' | tee -a /etc/apt/sources.list.d/llvm.list >/dev/null RUN apt-get install --yes zstd RUN apt-get install --yes libhwloc-dev diff --git a/.github/dockerfiles/Dockerfile_22.04-x86-64 b/.github/dockerfiles/Dockerfile_22.04-x86-64 index 0749cbaf26..9f6920226e 100644 --- a/.github/dockerfiles/Dockerfile_22.04-x86-64 +++ b/.github/dockerfiles/Dockerfile_22.04-x86-64 @@ -11,12 +11,11 @@ RUN dpkg --add-architecture riscv64 # The main archive only hosts amd64 and i386, we need to add ports for arm64 and riscv64. RUN sed -i -e '/^deb /{h;s|deb |&[arch=amd64,i386] |p;g;s|deb http://[^ ]*|deb [arch=arm64,riscv64] http://ports.ubuntu.com/ubuntu-ports|p;d}' /etc/apt/sources.list -# Add llvm repos -RUN wget -qO - https://apt.llvm.org/llvm-snapshot.gpg.key | gpg --dearmor - | tee /usr/share/keyrings/llvm-archive-keyring.gpg >/dev/null -RUN echo 'deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/jammy/ llvm-toolchain-jammy-19 main' | tee /etc/apt/sources.list.d/llvm.list >/dev/null -RUN echo 'deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/jammy/ llvm-toolchain-jammy-20 main' | tee -a /etc/apt/sources.list.d/llvm.list >/dev/null RUN apt-get update +# Add llvm key +RUN wget -qO - https://apt.llvm.org/llvm-snapshot.gpg.key | gpg --dearmor - | tee /usr/share/keyrings/llvm-archive-keyring.gpg >/dev/null + # Install minimum requirements RUN apt-get install --yes cmake libtinfo-dev # Install 32-bit requirements diff --git a/.github/dockerfiles/Dockerfile_24.04-x86-64 b/.github/dockerfiles/Dockerfile_24.04-x86-64 index 885932b22c..1cad2527c8 100644 --- a/.github/dockerfiles/Dockerfile_24.04-x86-64 +++ b/.github/dockerfiles/Dockerfile_24.04-x86-64 @@ -3,6 +3,7 @@ FROM ubuntu:24.04 RUN apt-get update RUN apt-get install -y wget git build-essential +# Enable ability to install foreign packages for cross compilation RUN dpkg --add-architecture i386 RUN dpkg --add-architecture arm64 RUN dpkg --add-architecture riscv64 @@ -12,10 +13,8 @@ RUN sed -i -e '/^Types:/,/^Signed-By:/{/Types:/{h;d};H;/^Signed-By:/{g;s|$|\nArc RUN apt-get update +# Add llvm key RUN wget -qO - https://apt.llvm.org/llvm-snapshot.gpg.key | gpg --dearmor - | tee /usr/share/keyrings/llvm-archive-keyring.gpg >/dev/null -RUN echo 'deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/noble/ llvm-toolchain-noble-19 main' | tee /etc/apt/sources.list.d/llvm.list >/dev/null -RUN echo 'deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/noble/ llvm-toolchain-noble-20 main' | tee -a /etc/apt/sources.list.d/llvm.list >/dev/null -RUN apt-get update # Install minimum requirements RUN apt-get install --yes cmake libtinfo-dev diff --git a/.github/workflows/codeql.yml b/.github/workflows/codeql.yml index 8385d1f7e4..d8279f1409 100644 --- a/.github/workflows/codeql.yml +++ b/.github/workflows/codeql.yml @@ -25,11 +25,11 @@ jobs: - name: Checkout repo uses: actions/checkout@v4 - # installs tools, ninja and installs llvm (default 19, RelAssert) and sets up cache + # installs tools, ninja and installs llvm (default 20, RelAssert) and sets up cache - name: setup-ubuntu uses: ./.github/actions/setup_build with: - llvm_version: 19 + llvm_version: 20 llvm_build_type: RelAssert # Initializes the CodeQL tools for scanning. @@ -81,7 +81,7 @@ jobs: - name: setup-ubuntu uses: ./.github/actions/setup_build with: - llvm_version: 19 + llvm_version: 20 llvm_build_type: RelAssert # Initializes the CodeQL tools for scanning. diff --git a/.github/workflows/create_llvm_artefacts.yml b/.github/workflows/create_llvm_artefacts.yml index 5938f6a074..715ecb99da 100644 --- a/.github/workflows/create_llvm_artefacts.yml +++ b/.github/workflows/create_llvm_artefacts.yml @@ -9,7 +9,7 @@ on: llvm_version: required: true type: string - description: 'llvm (major) version - 19,20 or main' + description: 'llvm (major) version - 20, 21 or main' use_github_cache: required: false type: boolean diff --git a/.github/workflows/create_publish_artifacts.yml b/.github/workflows/create_publish_artifacts.yml index dc80618e21..ef3999a4e6 100644 --- a/.github/workflows/create_publish_artifacts.yml +++ b/.github/workflows/create_publish_artifacts.yml @@ -33,7 +33,7 @@ jobs: - name: Setup ubuntu uses: ./.github/actions/setup_build with: - llvm_version: 19 + llvm_version: 20 llvm_build_type: Release - name: Setup python diff --git a/.github/workflows/planned_testing_caller.yml b/.github/workflows/planned_testing_caller.yml index 9aca101bee..a40b1dc411 100644 --- a/.github/workflows/planned_testing_caller.yml +++ b/.github/workflows/planned_testing_caller.yml @@ -4,10 +4,10 @@ on: workflow_call: inputs: llvm_version: - description: 'llvm major version (e.g 19,20, main) - to be used for llvm specific fails' + description: 'llvm major version (e.g 20,21,main) - to be used for llvm specific fails' type: string llvm_branch: - description: 'The actual llvm branch to check out e.g. release/19.x' + description: 'The actual llvm branch to check out e.g. release/20.x' type: string ock: type: boolean diff --git a/.github/workflows/planned_testing_caller_19.yml b/.github/workflows/planned_testing_caller_19.yml deleted file mode 100644 index eb1f86499b..0000000000 --- a/.github/workflows/planned_testing_caller_19.yml +++ /dev/null @@ -1,71 +0,0 @@ -name: run planned tests for llvm 19 -on: - pull_request: - paths: - - '.github/workflows/planned_testing_caller_19.yml' - schedule: - # Run Sun, Wed at 7pm - - cron: '0 19 * * 0,3' - workflow_dispatch: - inputs: - dpcpp_repo_pr_no: - description: 'dpc++ repo ref: enter PR number (else default branch is used)' - type: string - default: "" - -permissions: - packages: read - -concurrency: - group: ${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }} - cancel-in-progress: true - -jobs: - L19_: - if: github.repository == 'uxlfoundation/oneapi-construction-kit' || github.event_name == 'workflow_dispatch' - uses: ./.github/workflows/planned_testing_caller.yml - with: - llvm_version: '19' - llvm_branch: 'release/19.x' - use_llvm_github_cache: true - save_cache: ${{ github.event_name == 'schedule' }} - dpcpp_repo_pr_no: ${{ inputs.dpcpp_repo_pr_no }} - # We can set ock, test_sycl_cts etc here optionally if this is used as a pull request. - # Any with parameters below this is intended for local testing and should not be merged. - # target_list: '[ "host_x86_64_linux" ]' - # ock: true - # native_cpu: true - # test_tornado: true - # test_remote_hal: true - # test_sycl_cts: true - # test_sycl_e2e: true - # test_opencl_cts: true - # test_sanitizers: true - # run_internal: true - # run_external: true - # build_llvm: true - - # # These can be set to a workflow run-id to download from a previous workflow. This can be seen from - # # the action run of a job or workflow - # # e.g. https://github.com/uxlfoundation/oneapi-construction-kit/actions/runs//..... - # download_ock_artefact: host_x86_64_linux=14225268091;host_aarch64_linux=14225268091;host_riscv64_linux=14225268091 - # download_dpcpp_artefact: host_x86_64_linux=14225268091;host_aarch64_linux=14225268091;host_riscv64_linux=14225268091 - # download_sycl_cts_artefact: host_x86_64_linux=14225268091;host_aarch64_linux=14225268091;host_riscv64_linux=14225268091 - - # This cleans up any caches which may have been created when running external tests - clean_caches: - needs: [ L19_ ] - if: ( github.repository == 'uxlfoundation/oneapi-construction-kit' || github.event_name == 'workflow_dispatch' ) && always() - runs-on: ubuntu-latest - permissions: - actions: write # for gh cache delete - steps: - - name: Checkout repo - uses: actions/checkout@v4 - with: - sparse-checkout: .github - - name: Cache clean - uses: ./.github/actions/clean_cache - with: - token: ${{ secrets.GITHUB_TOKEN }} - cache_prefixes: "ccache-ccache-dpcpp-build-host_x86_64_linux ccache-ccache-dpcpp-build-host_aarch64_linux ccache-ccache-dpcpp-build-host_riscv64_linux" diff --git a/.github/workflows/planned_testing_caller_20.yml b/.github/workflows/planned_testing_caller_20.yml index 79f2a975fa..11bef3f8e8 100644 --- a/.github/workflows/planned_testing_caller_20.yml +++ b/.github/workflows/planned_testing_caller_20.yml @@ -4,8 +4,8 @@ on: paths: - '.github/workflows/planned_testing_caller_20.yml' schedule: - # Run Mon, Thu at 7pm - - cron: '0 19 * * 1,4' + # Run Sun, Wed at 7pm + - cron: '0 19 * * 0,3' workflow_dispatch: inputs: dpcpp_repo_pr_no: diff --git a/.github/workflows/planned_testing_caller_21.yml b/.github/workflows/planned_testing_caller_21.yml index b9c43d918f..dffb0b9d34 100644 --- a/.github/workflows/planned_testing_caller_21.yml +++ b/.github/workflows/planned_testing_caller_21.yml @@ -4,8 +4,8 @@ on: paths: - '.github/workflows/planned_testing_caller_21.yml' schedule: - # Run Tue, Fri at 7pm - - cron: '0 19 * * 2,5' + # Run Mon, Thu at 7pm + - cron: '0 19 * * 1,4' workflow_dispatch: inputs: dpcpp_repo_pr_no: diff --git a/.github/workflows/planned_testing_caller_main.yml b/.github/workflows/planned_testing_caller_main.yml index e26001d2e2..95e2488e12 100644 --- a/.github/workflows/planned_testing_caller_main.yml +++ b/.github/workflows/planned_testing_caller_main.yml @@ -4,8 +4,8 @@ on: paths: - '.github/workflows/planned_testing_caller_main.yml' schedule: - # Run Sat at 7pm - - cron: '0 19 * * 6' + # Run Wed, Sat, Sun at 7pm + - cron: '0 19 * * 2,5,6' permissions: packages: read diff --git a/.github/workflows/pr_tests_cache.yml b/.github/workflows/pr_tests_cache.yml index 9aba4f4003..8978af9f11 100644 --- a/.github/workflows/pr_tests_cache.yml +++ b/.github/workflows/pr_tests_cache.yml @@ -21,8 +21,8 @@ on: workflow_dispatch: env: - llvm_previous: '19' - llvm_current: '20' + llvm_previous: '20' + llvm_current: '21' concurrency: group: pr-test-cache-${{ github.workflow }}-${{ github.event.pull_request.number || github.ref }} diff --git a/.github/workflows/run_ock_external_tests.yml b/.github/workflows/run_ock_external_tests.yml index 088e4bbee9..d192b93083 100644 --- a/.github/workflows/run_ock_external_tests.yml +++ b/.github/workflows/run_ock_external_tests.yml @@ -50,13 +50,13 @@ on: llvm_version: required: false type: string - description: 'llvm major version (e.g 19,20, main) - to be used for llvm specific fails' - default: 19 + description: 'llvm major version (e.g 20,21,main) - to be used for llvm specific fails' + default: 20 llvm_branch: required: false type: string - description: 'The actual llvm branch to check out e.g. release/19.x' - default: 'release/19.x' + description: 'The actual llvm branch to check out e.g. release/20.x' + default: 'release/20.x' download_ock_artefact: required: false type: string diff --git a/.github/workflows/run_ock_internal_tests.yml b/.github/workflows/run_ock_internal_tests.yml index af3dfe885e..50a4aef9ea 100644 --- a/.github/workflows/run_ock_internal_tests.yml +++ b/.github/workflows/run_ock_internal_tests.yml @@ -131,6 +131,7 @@ jobs: # Install requirements - name: Install clang-tidy and parallel run: | + echo "deb [signed-by=/usr/share/keyrings/llvm-archive-keyring.gpg] http://apt.llvm.org/jammy/ llvm-toolchain-jammy-20 main" | sudo tee /etc/apt/sources.list.d/llvm.list >/dev/null sudo apt-get update sudo apt-get install --yes clang-tidy-20 parallel diff --git a/.github/workflows/run_pr_tests_caller.yml b/.github/workflows/run_pr_tests_caller.yml index 3ab9c9c724..24dd3320e7 100644 --- a/.github/workflows/run_pr_tests_caller.yml +++ b/.github/workflows/run_pr_tests_caller.yml @@ -30,5 +30,5 @@ jobs: with: is_pull_request: true llvm_source: install - llvm_previous: 19 - llvm_current: 20 + llvm_previous: 20 + llvm_current: 21 diff --git a/CHANGELOG.md b/CHANGELOG.md index 68c97fe3e9..87f4a65327 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -4,13 +4,13 @@ Feature additions: -* Support for LLVM 20 has been added. +* Support for LLVM 20 and 21 has been added. Upgrade guidance: * Support for degenerate subgroups has been removed. No in-tree target or template was using this, but custom targets may need to be updated. -* Support for LLVM versions before 19 has been removed. +* Support for LLVM versions before 20 has been removed. * Vulkan API support has been removed and will no longer be supported for any versions. Users may fork from the 4.0.0 version if they wish to use it. * The CMake option CA_ENABLE_API has been removed; OpenCL will always be enabled diff --git a/README.md b/README.md index 78daea55ae..2070db58ce 100644 --- a/README.md +++ b/README.md @@ -1,8 +1,8 @@ [![Contributor Covenant](https://img.shields.io/badge/Contributor%20Covenant-1.4-4baaaa.svg)](CODE_OF_CONDUCT.md) [![OpenSSF Scorecard](https://api.scorecard.dev/projects/github.com/uxlfoundation/oneapi-construction-kit/badge)](https://scorecard.dev/viewer/?uri=github.com/uxlfoundation/oneapi-construction-kit) -![planned testing LLVM 19](https://github.com/uxlfoundation/oneapi-construction-kit/actions/workflows/planned_testing_caller_19.yml/badge.svg) ![planned testing LLVM 20](https://github.com/uxlfoundation/oneapi-construction-kit/actions/workflows/planned_testing_caller_20.yml/badge.svg) +![planned testing LLVM 21](https://github.com/uxlfoundation/oneapi-construction-kit/actions/workflows/planned_testing_caller_21.yml/badge.svg) ![planned experimental testing LLVM main](https://github.com/uxlfoundation/oneapi-construction-kit/actions/workflows/planned_testing_caller_main.yml/badge.svg) [![UXL Foundation Logo](https://github.com/uxlfoundation/artwork/blob/main/foundation/uxl-foundation-logo-horizontal-color.png)][UXL Foundation] diff --git a/cmake/ImportLLVM.cmake b/cmake/ImportLLVM.cmake index 72c9d4d92b..36102278ee 100644 --- a/cmake/ImportLLVM.cmake +++ b/cmake/ImportLLVM.cmake @@ -57,8 +57,8 @@ include(DetectLLVMMSVCCRT) :cmake:variable:`LLVM_PACKAGE_VERSION` from the imported LLVM install is a supported version. #]=======================================================================] -set(CA_LLVM_MINIMUM_VERSION 19.0.0) -set(CA_LLVM_MAXIMUM_VERSION 20) +set(CA_LLVM_MINIMUM_VERSION 20.0.0) +set(CA_LLVM_MAXIMUM_VERSION 21) string(REPLACE ";" "', '" CA_LLVM_VERSIONS_PRETTY "${CA_LLVM_VERSIONS}") if("${LLVM_PACKAGE_VERSION}" VERSION_LESS "${CA_LLVM_MINIMUM_VERSION}") message(FATAL_ERROR diff --git a/doc/overview/compiler/supported-llvm-versions.rst b/doc/overview/compiler/supported-llvm-versions.rst index d9cc4ad7e0..97592b63e2 100644 --- a/doc/overview/compiler/supported-llvm-versions.rst +++ b/doc/overview/compiler/supported-llvm-versions.rst @@ -8,5 +8,5 @@ All targets .. rubric:: Supported - tested regularly -- 19.1.7+ -- 20.1.2+ +- 20.1.8+ +- 21.1.0+ diff --git a/modules/compiler/compiler_pipeline/CMakeLists.txt b/modules/compiler/compiler_pipeline/CMakeLists.txt index 94af9caf47..826485ad3b 100644 --- a/modules/compiler/compiler_pipeline/CMakeLists.txt +++ b/modules/compiler/compiler_pipeline/CMakeLists.txt @@ -48,7 +48,6 @@ add_ca_library(compiler-pipeline STATIC ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/pipeline_parse_helpers.h ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/prepare_barriers_pass.h ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/reduce_to_function_pass.h - ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/remove_address_spaces_pass.h ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/remove_exceptions_pass.h ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/remove_lifetime_intrinsics_pass.h ${CMAKE_CURRENT_SOURCE_DIR}/include/compiler/utils/rename_builtins_pass.h @@ -97,7 +96,6 @@ add_ca_library(compiler-pipeline STATIC ${CMAKE_CURRENT_SOURCE_DIR}/source/pass_machinery.cpp ${CMAKE_CURRENT_SOURCE_DIR}/source/prepare_barriers_pass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/source/reduce_to_function_pass.cpp - ${CMAKE_CURRENT_SOURCE_DIR}/source/remove_address_spaces_pass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/source/remove_exceptions_pass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/source/remove_lifetime_intrinsics_pass.cpp ${CMAKE_CURRENT_SOURCE_DIR}/source/rename_builtins_pass.cpp diff --git a/modules/compiler/compiler_pipeline/include/compiler/utils/remove_address_spaces_pass.h b/modules/compiler/compiler_pipeline/include/compiler/utils/remove_address_spaces_pass.h deleted file mode 100644 index 3f915381e2..0000000000 --- a/modules/compiler/compiler_pipeline/include/compiler/utils/remove_address_spaces_pass.h +++ /dev/null @@ -1,43 +0,0 @@ -// Copyright (C) Codeplay Software Limited -// -// Licensed under the Apache License, Version 2.0 (the "License") with LLVM -// Exceptions; you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT -// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the -// License for the specific language governing permissions and limitations -// under the License. -// -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -/// @file -/// -/// Remove address spaces pass. - -#ifndef COMPILER_UTILS_REMOVE_ADDRESS_SPACES_PASS_H_INCLUDED -#define COMPILER_UTILS_REMOVE_ADDRESS_SPACES_PASS_H_INCLUDED - -#include -#include - -#if LLVM_VERSION_LESS(20, 0) -namespace compiler { -namespace utils { - -/// @brief A pass that removes address spaces from functions to work around LLVM -/// issue #124759. -class RemoveAddressSpacesPass final - : public llvm::PassInfoMixin { - public: - llvm::PreservedAnalyses run(llvm::Function &, - llvm::FunctionAnalysisManager &); -}; -} // namespace utils -} // namespace compiler -#endif - -#endif // COMPILER_UTILS_REMOVE_ADDRESS_SPACES_PASS_H_INCLUDED diff --git a/modules/compiler/compiler_pipeline/source/builtin_info.cpp b/modules/compiler/compiler_pipeline/source/builtin_info.cpp index 3242031d61..dc85cd0ad5 100644 --- a/modules/compiler/compiler_pipeline/source/builtin_info.cpp +++ b/modules/compiler/compiler_pipeline/source/builtin_info.cpp @@ -545,7 +545,7 @@ Function *BuiltinInfo::getScalarEquivalent(const Builtin &B, Module *M) { Type *ScalarType = VecRetTy->getElementType(); // Get the scalar version of the intrinsic Function *ScalarIntrinsic = - multi_llvm::GetOrInsertIntrinsicDeclaration(M, IntrinsicID, ScalarType); + Intrinsic::getOrInsertDeclaration(M, IntrinsicID, ScalarType); return ScalarIntrinsic; } diff --git a/modules/compiler/compiler_pipeline/source/optimal_builtin_replacement_pass.cpp b/modules/compiler/compiler_pipeline/source/optimal_builtin_replacement_pass.cpp index 8323892827..e7610f89c5 100644 --- a/modules/compiler/compiler_pipeline/source/optimal_builtin_replacement_pass.cpp +++ b/modules/compiler/compiler_pipeline/source/optimal_builtin_replacement_pass.cpp @@ -28,7 +28,6 @@ #include #include #include -#include #include #define DEBUG_TYPE "ca-optimal-builtins" @@ -63,7 +62,7 @@ Value *OptimalBuiltinReplacementPass::replaceAbacusCLZ( // Get the declaration for the intrinsic auto *const ArgTy = Args[0]->getType(); auto *const Intrinsic = - multi_llvm::GetOrInsertIntrinsicDeclaration(M, Intrinsic::ctlz, ArgTy); + llvm::Intrinsic::getOrInsertDeclaration(M, Intrinsic::ctlz, ArgTy); // If we didn't find the intrinsic or the return type isn't what we // expect, skip this optimization Function *Callee = CB.getCalledFunction(); diff --git a/modules/compiler/compiler_pipeline/source/remove_address_spaces_pass.cpp b/modules/compiler/compiler_pipeline/source/remove_address_spaces_pass.cpp deleted file mode 100644 index 2accf29655..0000000000 --- a/modules/compiler/compiler_pipeline/source/remove_address_spaces_pass.cpp +++ /dev/null @@ -1,152 +0,0 @@ -// Copyright (C) Codeplay Software Limited -// -// Licensed under the Apache License, Version 2.0 (the "License") with LLVM -// Exceptions; you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT -// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the -// License for the specific language governing permissions and limitations -// under the License. -// -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#include -#include -#include -#include - -using namespace llvm; - -#if LLVM_VERSION_LESS(20, 0) -static Type *removeAddressSpaces(Type *T) { - constexpr unsigned DefaultAddressSpace = 0; - - if (auto *PtrTy = dyn_cast(T)) { - if (PtrTy->getAddressSpace() != DefaultAddressSpace) { - return PointerType::get(T->getContext(), DefaultAddressSpace); - } - } else if (auto *VecTy = dyn_cast(T)) { - if (auto *EltTy = removeAddressSpaces(VecTy->getElementType())) { - return VectorType::get(EltTy, VecTy); - } - } - - return nullptr; -} - -PreservedAnalyses compiler::utils::RemoveAddressSpacesPass::run( - Function &F, FunctionAnalysisManager &) { - bool Changed = false; - - auto MaybeCastOperand = [&](Instruction &I, Use &Op, Type *T) { - if (Op->getType() == T) return false; - IRBuilder<> B(&I); - Op.set(B.CreateAddrSpaceCast(Op, T)); - return true; - }; - - auto MaybeCastResult = [&](Value &V, Type *OldTy, Type *NewTy, - BasicBlock::iterator InsertPt) { - if (V.use_empty()) return false; - auto *CastV = - new AddrSpaceCastInst(PoisonValue::get(OldTy), NewTy, "", InsertPt); - V.mutateType(NewTy); - V.replaceAllUsesWith(CastV); - V.mutateType(OldTy); - CastV->setOperand(0, &V); - return true; - }; - - // If any args have a non-default address space, replace their uses with a - // cast to the default address space inserted in the entry block. - auto InsertPt = F.getEntryBlock().getFirstNonPHIOrDbgOrAlloca(); - for (auto &Arg : F.args()) { - if (auto *T = Arg.getType(); auto *NewTy = removeAddressSpaces(T)) { - Changed |= MaybeCastResult(Arg, T, NewTy, InsertPt); - } - } - - // Mutate all instructions to remove non-default address spaces. In most cases - // this is done by mutating the instruction directly, but in calls and - // extractvalues we cannot do that so insert a new addrspacecast instruction - // instead. - for (auto &BB : F) { - for (auto &I : BB) { - if (auto *T = I.getType(); auto *NewTy = removeAddressSpaces(T)) { - // call and extractvalue instructions need special handling. Cast their - // result instead. - if (isa(I) || isa(I)) { - InsertPt = std::next(I.getIterator()); - InsertPt.setHeadBit(true); - Changed |= MaybeCastResult(I, T, NewTy, InsertPt); - } else { - I.mutateType(NewTy); - Changed = true; - } - } - } - } - - // Now go over instructions again to remove address space casts made - // redundant, and insert new address space casts as required. - for (auto &BB : F) { - for (auto &I : make_early_inc_range(BB)) { - if (isa(I)) { - if (I.getType() == I.getOperand(0)->getType()) { - I.replaceAllUsesWith(I.getOperand(0)); - I.eraseFromParent(); - Changed = true; - } - continue; - } - - // For call instructions, operand types need to match parameter types. - if (auto *CB = dyn_cast(&I)) { - auto *FTy = CB->getFunctionType(); - for (unsigned Idx = 0, E = FTy->getNumParams(); Idx != E; ++Idx) { - Changed |= - MaybeCastOperand(I, I.getOperandUse(Idx), FTy->getParamType(Idx)); - } - continue; - } - - // For ret instructions, operand types need to match return type. - if (auto *RI = dyn_cast(&I); RI && RI->getReturnValue()) { - Changed |= MaybeCastOperand(I, I.getOperandUse(0), F.getReturnType()); - continue; - } - - // For insertvalue instructions, operand types need to match structure or - // array element type. - if (auto *IVI = dyn_cast(&I)) { - Changed |= MaybeCastOperand( - I, I.getOperandUse(InsertValueInst::getInsertedValueOperandIndex()), - ExtractValueInst::getIndexedType( - IVI->getAggregateOperand()->getType(), IVI->getIndices())); - continue; - } - - // For other instructions, operands should be non-address-space-qualified. - // Operands that are arguments or other instructions will have been - // updated already, but we may still have constants to deal with. - for (auto &Op : I.operands()) { - if (auto *T = Op->getType(); auto *NewTy = removeAddressSpaces(T)) { - Changed |= MaybeCastOperand(I, Op, NewTy); - } - } - } - } - - if (Changed) { - PreservedAnalyses PA; - PA.preserveSet(); - return PA; - } else { - return PreservedAnalyses::all(); - } -} -#endif diff --git a/modules/compiler/multi_llvm/include/multi_llvm/dibuilder.h b/modules/compiler/multi_llvm/include/multi_llvm/dibuilder.h deleted file mode 100644 index 487315a807..0000000000 --- a/modules/compiler/multi_llvm/include/multi_llvm/dibuilder.h +++ /dev/null @@ -1,101 +0,0 @@ -// Copyright (C) Codeplay Software Limited -// -// Licensed under the Apache License, Version 2.0 (the "License") with LLVM -// Exceptions; you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT -// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the -// License for the specific language governing permissions and limitations -// under the License. -// -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef MULTI_LLVM_DIBUILDER_H_INCLUDED -#define MULTI_LLVM_DIBUILDER_H_INCLUDED - -#include -#include - -#include - -namespace multi_llvm { -// TODO In order to enable the use of OCK in DPC++ which currently uses the -// older DIBuilder interface, we do not yet condition this on LLVM version, we -// dynamically detect which version of DIBuilder we have. This should be updated -// after DPC++'s next pulldown to drop the use of DIBuilderWrapperNeeded and -// base it entirely on LLVM major version. -#if LLVM_VERSION_GREATER_EQUAL(20, 0) && 0 -using DIBuilder = llvm::DIBuilder; -#else -template -struct DIBuilderWrapper : DIBuilder { - using DIBuilder::DIBuilder; - - llvm::BasicBlock *getBasicBlock(llvm::InsertPosition InsertPt) { - return InsertPt.getBasicBlock(); - } - - auto insertDeclare(llvm::Value *Storage, llvm::DILocalVariable *VarInfo, - llvm::DIExpression *Expr, const llvm::DILocation *DL, - llvm::BasicBlock *InsertAtEnd) { - return DIBuilder::insertDeclare(Storage, VarInfo, Expr, DL, InsertAtEnd); - } - - auto insertDeclare(llvm::Value *Storage, llvm::DILocalVariable *VarInfo, - llvm::DIExpression *Expr, const llvm::DILocation *DL, - llvm::BasicBlock::iterator InsertPt) { - auto *InsertBB = getBasicBlock(InsertPt); - if (InsertPt == InsertBB->end()) { - return DIBuilder::insertDeclare(Storage, VarInfo, Expr, DL, InsertBB); - } else { - return DIBuilder::insertDeclare(Storage, VarInfo, Expr, DL, &*InsertPt); - } - } - - auto insertDbgValueIntrinsic(llvm::Value *Val, llvm::DILocalVariable *VarInfo, - llvm::DIExpression *Expr, - const llvm::DILocation *DL, - llvm::BasicBlock *InsertAtEnd) { - return DIBuilder::insertDbgValueIntrinsic(Val, VarInfo, Expr, DL, - InsertAtEnd); - } - - auto insertDbgValueIntrinsic(llvm::Value *Val, llvm::DILocalVariable *VarInfo, - llvm::DIExpression *Expr, - const llvm::DILocation *DL, - llvm::BasicBlock::iterator InsertPt) { - auto *InsertBB = getBasicBlock(InsertPt); - if (InsertPt == InsertBB->end()) { - return DIBuilder::insertDbgValueIntrinsic(Val, VarInfo, Expr, DL, - InsertBB); - } else { - return DIBuilder::insertDbgValueIntrinsic(Val, VarInfo, Expr, DL, - &*InsertPt); - } - } -}; - -template -static constexpr bool DIBuilderWrapperNeeded = true; - -template -static constexpr bool DIBuilderWrapperNeeded< - DIBuilder, std::void_t().insertLabel( - std::declval(), - std::declval(), - std::declval()))>> = false; - -template -using DIBuilderMaybeWrapped = - std::conditional_t, - DIBuilderWrapper, DIBuilder>; - -using DIBuilder = DIBuilderMaybeWrapped; -#endif -} // namespace multi_llvm - -#endif // MULTI_LLVM_DIBUILDER_H_INCLUDED diff --git a/modules/compiler/multi_llvm/include/multi_llvm/instructions.h b/modules/compiler/multi_llvm/include/multi_llvm/instructions.h index 19ec81efc5..fabe42ae57 100644 --- a/modules/compiler/multi_llvm/include/multi_llvm/instructions.h +++ b/modules/compiler/multi_llvm/include/multi_llvm/instructions.h @@ -36,9 +36,6 @@ struct BinOpHelper; #define LLVM 20 #include #undef LLVM -#define LLVM 19 -#include -#undef LLVM } // namespace detail diff --git a/modules/compiler/multi_llvm/include/multi_llvm/instructions.inc b/modules/compiler/multi_llvm/include/multi_llvm/instructions.inc index 80ae42c23f..d9049b2584 100644 --- a/modules/compiler/multi_llvm/include/multi_llvm/instructions.inc +++ b/modules/compiler/multi_llvm/include/multi_llvm/instructions.inc @@ -14,20 +14,13 @@ // // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -#if LLVM == 19 -template -struct BinOpHelper> -#define BINOP_LLVM20(OP, STR) -#define BINOP_LLVM21(OP, STR) -#elif LLVM == 20 +#if LLVM == 20 template struct BinOpHelper> -#define BINOP_LLVM20(OP, STR) BINOP(OP, STR) #define BINOP_LLVM21(OP, STR) #elif LLVM == 21 template struct BinOpHelper -#define BINOP_LLVM20(OP, STR) BINOP(OP, STR) #define BINOP_LLVM21(OP, STR) BINOP(OP, STR) #endif { @@ -51,8 +44,8 @@ struct BinOpHelper BINOP_LLVM21(FMinimum, "fminumum") \ BINOP(UIncWrap, "uincwrap") \ BINOP(UDecWrap, "udecwrap") \ - BINOP_LLVM20(USubCond, "usubcond") \ - BINOP_LLVM20(USubSat, "usubsat") + BINOP(USubCond, "usubcond") \ + BINOP(USubSat, "usubsat") static std::optional consume_front_with_underscore( llvm::StringRef &String) { @@ -79,6 +72,5 @@ struct BinOpHelper } #undef BINOPS -#undef BINOP_LLVM20 #undef BINOP_LLVM21 }; diff --git a/modules/compiler/multi_llvm/include/multi_llvm/intrinsic.h b/modules/compiler/multi_llvm/include/multi_llvm/intrinsic.h index f8e5800fa9..3df2c19ae8 100644 --- a/modules/compiler/multi_llvm/include/multi_llvm/intrinsic.h +++ b/modules/compiler/multi_llvm/include/multi_llvm/intrinsic.h @@ -22,19 +22,7 @@ namespace multi_llvm { -static inline auto GetOrInsertIntrinsicDeclaration( - llvm::Module *M, llvm::Intrinsic::ID id, - llvm::ArrayRef Tys = {}) { -#if LLVM_VERSION_GREATER_EQUAL(20, 0) - return llvm::Intrinsic::getOrInsertDeclaration(M, id, Tys); -#else - return llvm::Intrinsic::getDeclaration(M, id, Tys); -#endif -} - -// Drop getAttributes workaround when LLVM 20 is minimum version -// This can also be simplified once DPC++ catches up with getAttributes -// with FunctionType as the last argument. +// Drop getAttributes workaround when LLVM 21 is minimum version namespace detail { template auto getAttributes(T... args) diff --git a/modules/compiler/multi_llvm/include/multi_llvm/loop_utils.h b/modules/compiler/multi_llvm/include/multi_llvm/loop_utils.h deleted file mode 100644 index 2d6e8ba84b..0000000000 --- a/modules/compiler/multi_llvm/include/multi_llvm/loop_utils.h +++ /dev/null @@ -1,37 +0,0 @@ -// Copyright (C) Codeplay Software Limited -// -// Licensed under the Apache License, Version 2.0 (the "License") with LLVM -// Exceptions; you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// https://github.com/uxlfoundation/oneapi-construction-kit/blob/main/LICENSE.txt -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, WITHOUT -// WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. See the -// License for the specific language governing permissions and limitations -// under the License. -// -// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception - -#ifndef MULTI_LLVM_LOOP_UTILS_H_INCLUDED -#define MULTI_LLVM_LOOP_UTILS_H_INCLUDED - -#include -#include - -namespace multi_llvm { - -inline llvm::Value *createSimpleReduction(llvm::IRBuilderBase &B, - llvm::Value *Src, - llvm::RecurKind RdxKind) { -#if LLVM_VERSION_MAJOR >= 20 - return llvm::createSimpleReduction(B, Src, RdxKind); -#else - return llvm::createSimpleTargetReduction(B, Src, RdxKind); -#endif -} - -} // namespace multi_llvm - -#endif // MULTI_LLVM_LOOP_UTILS_H_INCLUDED diff --git a/modules/compiler/multi_llvm/include/multi_llvm/multi_llvm.h b/modules/compiler/multi_llvm/include/multi_llvm/multi_llvm.h index f80557fcf4..34ee670744 100644 --- a/modules/compiler/multi_llvm/include/multi_llvm/multi_llvm.h +++ b/modules/compiler/multi_llvm/include/multi_llvm/multi_llvm.h @@ -18,6 +18,5 @@ #define MULTI_LLVM_MULTI_LLVM_H_INCLUDED #include -#include #endif // MULTI_LLVM_MULTI_LLVM_H_INCLUDED diff --git a/modules/compiler/source/base/include/base/target.h b/modules/compiler/source/base/include/base/target.h index 1865bdfe4a..32f3a65c48 100644 --- a/modules/compiler/source/base/include/base/target.h +++ b/modules/compiler/source/base/include/base/target.h @@ -24,6 +24,7 @@ #include #include +#include #include namespace compiler { diff --git a/modules/compiler/source/base/source/base_module_pass_machinery.cpp b/modules/compiler/source/base/source/base_module_pass_machinery.cpp index 116f0b5270..cb5f97e3b9 100644 --- a/modules/compiler/source/base/source/base_module_pass_machinery.cpp +++ b/modules/compiler/source/base/source/base_module_pass_machinery.cpp @@ -47,7 +47,6 @@ #include #include #include -#include #include #include #include diff --git a/modules/compiler/source/base/source/base_module_pass_registry.def b/modules/compiler/source/base/source/base_module_pass_registry.def index 2bcd19cfb5..f96bc1e70c 100644 --- a/modules/compiler/source/base/source/base_module_pass_registry.def +++ b/modules/compiler/source/base/source/base_module_pass_registry.def @@ -175,9 +175,6 @@ FUNCTION_PASS("manual-type-legalization", compiler::utils::ManualTypeLegalizatio FUNCTION_PASS("software-div", compiler::SoftwareDivisionPass()) FUNCTION_PASS("replace-addrspace-fns", compiler::utils::ReplaceAddressSpaceQualifierFunctionsPass()) FUNCTION_PASS("remove-lifetime", compiler::utils::RemoveLifetimeIntrinsicsPass()) -#if LLVM_VERSION_LESS(20, 0) -FUNCTION_PASS("remove-address-spaces", compiler::utils::RemoveAddressSpacesPass()) -#endif FUNCTION_PASS("remove-exceptions", compiler::utils::RemoveExceptionsPass()) FUNCTION_PASS("replace-mem-intrins", compiler::utils::ReplaceMemIntrinsicsPass()) FUNCTION_PASS("print", compiler::utils::GenericMetadataPrinterPass(llvm::dbgs())) diff --git a/modules/compiler/source/base/source/module.cpp b/modules/compiler/source/base/source/module.cpp index 5b3ba915fc..269eee30ed 100644 --- a/modules/compiler/source/base/source/module.cpp +++ b/modules/compiler/source/base/source/module.cpp @@ -1474,9 +1474,7 @@ std::unique_ptr BaseModule::compileOpenCLCToIR( #endif // CA_ENABLE_DEBUG_SUPPORT instance.createDiagnostics( -#if LLVM_VERSION_GREATER_EQUAL(20, 0) *llvm::vfs::getRealFileSystem(), -#endif new FrontendDiagnosticPrinter(*this, instance.getDiagnosticOpts())); // Write a copy of the kernel source out to disk and update the debug info diff --git a/modules/compiler/spirv-ll/include/spirv-ll/builder.h b/modules/compiler/spirv-ll/include/spirv-ll/builder.h index 8b6a64a218..7fa41b8582 100644 --- a/modules/compiler/spirv-ll/include/spirv-ll/builder.h +++ b/modules/compiler/spirv-ll/include/spirv-ll/builder.h @@ -23,7 +23,6 @@ #include #include #include -#include #include #include #include @@ -816,7 +815,7 @@ class Builder { /// @brief The IRBuilder used to generate the LLVM IR llvm::IRBuilder<> IRBuilder; /// @brief The DIBuilder used to generate the LLVM IR debug instructions - multi_llvm::DIBuilder DIBuilder; + llvm::DIBuilder DIBuilder; /// @brief Function the builder is currently working on llvm::Function *CurrentFunction; /// @brief Kernel function the builder is currently working on, or nullptr if diff --git a/modules/compiler/spirv-ll/include/spirv-ll/builder_debug_info.h b/modules/compiler/spirv-ll/include/spirv-ll/builder_debug_info.h index 8026838895..4536e6abc2 100644 --- a/modules/compiler/spirv-ll/include/spirv-ll/builder_debug_info.h +++ b/modules/compiler/spirv-ll/include/spirv-ll/builder_debug_info.h @@ -53,7 +53,7 @@ class DebugInfoBuilder : public ExtInstSetHandler { uint64_t workarounds = 0; /// @brief Map from DebugInfo instructions to the llvm::DIBuilder that builds /// them. - std::unordered_map> + std::unordered_map> debug_builder_map; /// @brief Cache of translated DebugInfo instructions. @@ -74,13 +74,13 @@ class DebugInfoBuilder : public ExtInstSetHandler { /// We may only have one DICompileUnit per DIBuilder, so must support /// multiple builders. This function finds the DIBuilder for the instruction /// based on its chain of scopes, if applicable. - multi_llvm::DIBuilder &getDIBuilder(const OpExtInst *op) const; + llvm::DIBuilder &getDIBuilder(const OpExtInst *op) const; /// @brief Returns the first registered DIBuilder, for when it doesn't matter /// which is used. /// /// @see getDIBuilder - multi_llvm::DIBuilder &getDefaultDIBuilder() const; + llvm::DIBuilder &getDefaultDIBuilder() const; /// @brief Returns true if the given ID is DebugInfoNone. bool isDebugInfoNone(spv::Id id) const; diff --git a/modules/compiler/spirv-ll/source/builder_debug_info.cpp b/modules/compiler/spirv-ll/source/builder_debug_info.cpp index afc1af7669..c1b54f26dc 100644 --- a/modules/compiler/spirv-ll/source/builder_debug_info.cpp +++ b/modules/compiler/spirv-ll/source/builder_debug_info.cpp @@ -301,7 +301,7 @@ bool DebugInfoBuilder::isDebugInfoSet(uint32_t set_id) const { set == ExtendedInstrSet::OpenCLDebugInfo100; } -multi_llvm::DIBuilder &DebugInfoBuilder::getDefaultDIBuilder() const { +llvm::DIBuilder &DebugInfoBuilder::getDefaultDIBuilder() const { assert(debug_builder_map.size() != 0 && "No DIBuilders"); return *debug_builder_map.begin()->second; } @@ -527,12 +527,8 @@ llvm::Expected DebugInfoBuilder::translate( } if (flags & OpenCLDebugInfo100FlagObjectPointer) { -#if LLVM_VERSION_GREATER_EQUAL(20, 0) const bool implicit = flags & OpenCLDebugInfo100FlagArtificial; type = dib.createObjectPointerType(type, implicit); -#else - type = dib.createObjectPointerType(type); -#endif } else if (flags & OpenCLDebugInfo100FlagArtificial) { type = dib.createArtificialType(type); } @@ -2341,7 +2337,7 @@ llvm::Error DebugInfoBuilder::create(const OpExtInst &opc) { CREATE_CASE(OpenCLDebugInfo100DebugImportedEntity, DebugImportedEntity) case OpenCLDebugInfo100Instructions::OpenCLDebugInfo100DebugCompilationUnit: debug_builder_map[opc.IdResult()] = - std::make_unique(*module.llvmModule); + std::make_unique(*module.llvmModule); break; case OpenCLDebugInfo100Instructions::OpenCLDebugInfo100DebugFunction: { // Translate and register the DISubprogram for the function. @@ -2381,10 +2377,9 @@ llvm::Error DebugInfoBuilder::create(const OpExtInst &opc) { return llvm::Error::success(); } -multi_llvm::DIBuilder &DebugInfoBuilder::getDIBuilder( - const OpExtInst *op) const { +llvm::DIBuilder &DebugInfoBuilder::getDIBuilder(const OpExtInst *op) const { assert(debug_builder_map.size() != 0 && "No DIBuilders"); - multi_llvm::DIBuilder &default_dib = getDefaultDIBuilder(); + llvm::DIBuilder &default_dib = getDefaultDIBuilder(); assert(isDebugInfoSet(op->Set()) && "Unexpected extended instruction set"); diff --git a/modules/compiler/targets/host/source/AddFloatingPointControl.cpp b/modules/compiler/targets/host/source/AddFloatingPointControl.cpp index c46387dc78..5db6d9cedd 100644 --- a/modules/compiler/targets/host/source/AddFloatingPointControl.cpp +++ b/modules/compiler/targets/host/source/AddFloatingPointControl.cpp @@ -27,7 +27,6 @@ #include #include #include -#include using namespace llvm; @@ -43,11 +42,11 @@ void configX86FP(Function &wrapper, Function &function) { // x86 STMXCSR instruction stores the contents of the MXCSR register in the // destination operand. MXCSR contains flags for control and status // information regarding SSE instructions. - Function *const st_mxcsr = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *const st_mxcsr = Intrinsic::getOrInsertDeclaration( wrapper.getParent(), Intrinsic::x86_sse_stmxcsr); // Loads the source operand into the MXCSR register - Function *const ld_mxcsr = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *const ld_mxcsr = Intrinsic::getOrInsertDeclaration( wrapper.getParent(), Intrinsic::x86_sse_ldmxcsr); // Allocas to store the old and new state @@ -132,10 +131,10 @@ void configArmFP(Function &wrapper, Function &function) { // create an IR builder with a single basic block in our wrapper IRBuilder<> ir(BasicBlock::Create(wrapper.getContext(), "", &wrapper)); - Function *const get_fpscr = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *const get_fpscr = Intrinsic::getOrInsertDeclaration( wrapper.getParent(), Intrinsic::arm_get_fpscr); - Function *const set_fpscr = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *const set_fpscr = Intrinsic::getOrInsertDeclaration( wrapper.getParent(), Intrinsic::arm_set_fpscr); // call fpscr to store the original state diff --git a/modules/compiler/targets/host/source/HostPassMachinery.cpp b/modules/compiler/targets/host/source/HostPassMachinery.cpp index 74ec16d125..7ce26b6e60 100644 --- a/modules/compiler/targets/host/source/HostPassMachinery.cpp +++ b/modules/compiler/targets/host/source/HostPassMachinery.cpp @@ -31,7 +31,6 @@ #include #include #include -#include #include #include #include @@ -286,11 +285,6 @@ llvm::ModulePassManager HostPassMachinery::getKernelFinalizationPasses( // be inlined even at -O0. PM.addPass(llvm::AlwaysInlinerPass()); -#if LLVM_VERSION_LESS(20, 0) - PM.addPass(llvm::createModuleToFunctionPassAdaptor( - compiler::utils::RemoveAddressSpacesPass())); -#endif - // Running this pass here is the "nuclear option", it would be better to // ensure exception handling is never introduced in the first place, but // it is not always plausible to do. diff --git a/modules/compiler/test/lit/passes/barriers-dbg.ll b/modules/compiler/test/lit/passes/barriers-dbg.ll index af589883b8..371437fa80 100644 --- a/modules/compiler/test/lit/passes/barriers-dbg.ll +++ b/modules/compiler/test/lit/passes/barriers-dbg.ll @@ -14,8 +14,7 @@ ; ; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -; RUN: %pp-llvm-ver -o %t < %s --llvm-ver %LLVMVER -; RUN: muxc --passes "work-item-loops,verify" -S %s | FileCheck %t +; RUN: muxc --passes "work-item-loops,verify" -S %s | FileCheck %s target triple = "spir64-unknown-unknown" target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" @@ -134,28 +133,22 @@ if.end: ; preds = %if.then, %entry ; ; The pass non-determinism also affects the placement of padding elements which ; further complicates the issue since this too impacts the offsets. See CA-119. -; CHECK-GE20: #dbg_value(ptr poison -; CHECK-EQ19: #dbg_value(ptr undef +; CHECK: #dbg_value(ptr poison ; CHECK-SAME: !DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{0|[1-9][0-9]*}} ; -; CHECK-GE20: #dbg_value(ptr poison -; CHECK-EQ19: #dbg_value(ptr undef +; CHECK: #dbg_value(ptr poison ; CHECK-SAME: !DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{0|[1-9][0-9]*}} ; -; CHECK-GE20: #dbg_value(ptr poison -; CHECK-EQ19: #dbg_value(ptr undef +; CHECK: #dbg_value(ptr poison ; CHECK-SAME: !DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{0|[1-9][0-9]*}} ; -; CHECK-GE20: #dbg_value(ptr poison -; CHECK-EQ19: #dbg_value(ptr undef +; CHECK: #dbg_value(ptr poison ; CHECK-SAME: !DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{0|[1-9][0-9]*}} ; -; CHECK-GE20: #dbg_value(ptr poison -; CHECK-EQ19: #dbg_value(ptr undef +; CHECK: #dbg_value(ptr poison ; CHECK-SAME: !DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{0|[1-9][0-9]*}} ; -; CHECK-GE20: #dbg_value(ptr poison -; CHECK-EQ19: #dbg_value(ptr undef +; CHECK: #dbg_value(ptr poison ; CHECK-SAME: !DIExpression(DW_OP_deref, DW_OP_plus_uconst, {{0|[1-9][0-9]*}} ; ; Debug info for first kernel scope diff --git a/modules/compiler/utils/source/lld_linker.cpp b/modules/compiler/utils/source/lld_linker.cpp index aeed288daa..472d6ff716 100644 --- a/modules/compiler/utils/source/lld_linker.cpp +++ b/modules/compiler/utils/source/lld_linker.cpp @@ -148,18 +148,6 @@ Expected> lldLinkToBinary( } #endif // NDEBUG -#if LLVM_VERSION_LESS(20, 0) - // LLVM's register allocator has issues with early-clobbers if subreg liveness - // is enabled. The InitUndef pass documents this and attempts to work around - // it, but prior to , the - // InitUndef pass would not work reliably when multiple functions were - // processed, because internal state from one function would be kept around - // when processing the next. As we have no good way of fixing the InitUndef - // pass in older LLVM versions, disable subreg liveness instead. - args.push_back("-mllvm"); - args.push_back("-enable-subreg-liveness=false"); -#endif - for (const auto &arg : additionalLinkArgs) { args.push_back(arg); } diff --git a/modules/compiler/vecz/source/transform/packetization_helpers.cpp b/modules/compiler/vecz/source/transform/packetization_helpers.cpp index bac6b19789..996881c8b1 100644 --- a/modules/compiler/vecz/source/transform/packetization_helpers.cpp +++ b/modules/compiler/vecz/source/transform/packetization_helpers.cpp @@ -30,7 +30,6 @@ #include #include #include -#include #include #include "debugging.h" @@ -241,7 +240,7 @@ Value *createMaybeVPReduction(IRBuilderBase &B, Value *Val, RecurKind Kind, assert(isa(Val->getType()) && "Must be vector type"); // If VL is null, it's not a vector-predicated reduction. if (!VL) { - return multi_llvm::createSimpleReduction(B, Val, Kind); + return createSimpleReduction(B, Val, Kind); } auto IntrinsicOp = Intrinsic::not_intrinsic; switch (Kind) { @@ -290,7 +289,7 @@ Value *createMaybeVPReduction(IRBuilderBase &B, Value *Val, RecurKind Kind, break; } - auto *const F = multi_llvm::GetOrInsertIntrinsicDeclaration( + auto *const F = Intrinsic::getOrInsertDeclaration( B.GetInsertBlock()->getModule(), IntrinsicOp, Val->getType()); assert(F && "Could not declare vector-predicated reduction intrinsic"); diff --git a/modules/compiler/vecz/source/transform/pre_linearize_pass.cpp b/modules/compiler/vecz/source/transform/pre_linearize_pass.cpp index 05913dbe3c..e82ec366c8 100644 --- a/modules/compiler/vecz/source/transform/pre_linearize_pass.cpp +++ b/modules/compiler/vecz/source/transform/pre_linearize_pass.cpp @@ -120,7 +120,7 @@ InstructionCost calculateBoolReductionCost(LLVMContext &context, Module *module, auto *F = Function::Create(new_fty, Function::InternalLinkage, "tmp", module); auto *BB = BasicBlock::Create(context, "reduce", F); IRBuilder<> B(BB); - multi_llvm::createSimpleReduction(B, &*F->arg_begin(), RecurKind::And); + createSimpleReduction(B, &*F->arg_begin(), RecurKind::And); const InstructionCost cost = calculateBlockCost(*BB, TTI); // We don't really need that function in the module anymore because it's diff --git a/modules/compiler/vecz/source/transform/scalarizer.cpp b/modules/compiler/vecz/source/transform/scalarizer.cpp index 5b79287c14..63b33c02f3 100644 --- a/modules/compiler/vecz/source/transform/scalarizer.cpp +++ b/modules/compiler/vecz/source/transform/scalarizer.cpp @@ -20,6 +20,7 @@ #include #include #include +#include #include #include #include @@ -27,7 +28,6 @@ #include #include #include -#include #include #include @@ -683,7 +683,7 @@ void Scalarizer::scalarizeDI(Instruction *Original, const SimdPacket *Packet, // instructions and is used to avoid duplicate LLVM dbg.value's. SmallPtrSet VectorElements; - multi_llvm::DIBuilder DIB(*Original->getModule(), false); + DIBuilder DIB(*Original->getModule(), false); for (DbgVariableRecord *const DVR : LAM->getAllDbgVariableRecordUsers()) { DILocalVariable *DILocal = nullptr; diff --git a/modules/compiler/vecz/source/vector_target_info.cpp b/modules/compiler/vecz/source/vector_target_info.cpp index e019994cb7..2ab49546fe 100644 --- a/modules/compiler/vecz/source/vector_target_info.cpp +++ b/modules/compiler/vecz/source/vector_target_info.cpp @@ -20,7 +20,6 @@ #include #include #include -#include #include #include @@ -504,7 +503,7 @@ Value *TargetInfo::createMaskedGatherLoad(IRBuilder<> &B, Type *Ty, Value *Ptr, const SmallVector Tys = {Ty, VecPtrTy}; return B.CreateIntrinsic(llvm::Intrinsic::vp_gather, Tys, Args); } else if (Legality.isMaskLegal()) { - Function *MaskedGather = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *MaskedGather = Intrinsic::getOrInsertDeclaration( F->getParent(), Intrinsic::masked_gather, {Ty, VecPtrTy}); if (MaskedGather) { @@ -606,7 +605,7 @@ Value *TargetInfo::createMaskedScatterStore(IRBuilder<> &B, Value *Data, const SmallVector Tys = {Data->getType(), VecPtrTy}; return B.CreateIntrinsic(llvm::Intrinsic::vp_scatter, Tys, Args); } else if (Legality.isMaskLegal()) { - Function *MaskedScatter = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *MaskedScatter = Intrinsic::getOrInsertDeclaration( F->getParent(), Intrinsic::masked_scatter, {DataTy, VecPtrTy}); if (MaskedScatter) { diff --git a/modules/compiler/vecz/source/vector_target_info_arm.cpp b/modules/compiler/vecz/source/vector_target_info_arm.cpp index 892c81483c..148db14475 100644 --- a/modules/compiler/vecz/source/vector_target_info_arm.cpp +++ b/modules/compiler/vecz/source/vector_target_info_arm.cpp @@ -18,7 +18,6 @@ #include #include #include -#include #include #include "debugging.h" @@ -224,7 +223,7 @@ bool TargetInfoArm::optimizeInterleavedGroup(IRBuilder<> &B, Tys = {VecTy, PtrTy}; } - Function *IntrFn = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *IntrFn = Intrinsic::getOrInsertDeclaration( Op0->getModule(), (Intrinsic::ID)IntrID, Tys); if (!IntrFn) { return false; @@ -379,7 +378,7 @@ bool TargetInfoAArch64::optimizeInterleavedGroup( VecTy = cast(Op0->getType()); } - Function *IntrFn = multi_llvm::GetOrInsertIntrinsicDeclaration( + Function *IntrFn = Intrinsic::getOrInsertDeclaration( Op0->getModule(), (Intrinsic::ID)IntrID, {VecTy, PtrTy}); if (!IntrFn) { return false; diff --git a/modules/compiler/vecz/source/vectorization_helpers.cpp b/modules/compiler/vecz/source/vectorization_helpers.cpp index 38e7752532..24c892f7b6 100644 --- a/modules/compiler/vecz/source/vectorization_helpers.cpp +++ b/modules/compiler/vecz/source/vectorization_helpers.cpp @@ -19,11 +19,11 @@ #include #include #include +#include #include #include #include #include -#include #include @@ -274,7 +274,7 @@ void cloneDebugInfo(const VectorizationUnit &VU) { } // Create a DISubprogram entry for the vectorized kernel - multi_llvm::DIBuilder DIB(*VU.scalarFunction()->getParent(), false); + DIBuilder DIB(*VU.scalarFunction()->getParent(), false); DICompileUnit *CU = DIB.createCompileUnit(dwarf::DW_LANG_OpenCL, ScalarDI->getFile(), "", ScalarDI->isOptimized(), "", 0); diff --git a/modules/compiler/vecz/test/lit/llvm/VectorPredication/packetize_mask_varying.ll b/modules/compiler/vecz/test/lit/llvm/VectorPredication/packetize_mask_varying.ll index 5c1df71ed9..0ce65b9f4c 100644 --- a/modules/compiler/vecz/test/lit/llvm/VectorPredication/packetize_mask_varying.ll +++ b/modules/compiler/vecz/test/lit/llvm/VectorPredication/packetize_mask_varying.ll @@ -14,8 +14,7 @@ ; ; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -; RUN: %pp-llvm-ver -o %t < %s --llvm-ver %LLVMVER -; RUN: veczc -k mask_varying -vecz-scalable -vecz-simd-width=4 -vecz-choices=VectorPredication -S < %s | FileCheck %t +; RUN: veczc -k mask_varying -vecz-scalable -vecz-simd-width=4 -vecz-choices=VectorPredication -S < %s | FileCheck %s target triple = "spir64-unknown-unknown" target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" @@ -40,10 +39,7 @@ if.end: ; CHECK: define spir_kernel void @__vecz_nxv4_vp_mask_varying ; CHECK: [[CMP:%.*]] = icmp slt %{{.*}}, ; CHECK: [[RED:%.*]] = call i1 @llvm.vp.reduce.or.nxv4i1(i1 false, [[CMP]], {{.*}}, i32 {{.*}}) -; CHECK-LT20: [[REINS:%.*]] = insertelement <4 x i1> poison, i1 [[RED]], {{(i32|i64)}} 0 -; CHECK-LT20: [[RESPLAT:%.*]] = shufflevector <4 x i1> [[REINS]], <4 x i1> poison, <4 x i32> zeroinitializer -; CHECK-LT20: [[VAL:%.*]] = call <4 x i32> @__vecz_b_masked_load16_Dv4_ju3ptrDv4_b(ptr %aptr, <4 x i1> [[RESPLAT]]) -; CHECK-GE20: [[VAL:%.*]] = load <4 x i32>, ptr %aptr +; CHECK: [[VAL:%.*]] = load <4 x i32>, ptr %aptr } declare i64 @__mux_get_global_id(i32) diff --git a/modules/compiler/vecz/test/lit/llvm/irreducible_loop.ll b/modules/compiler/vecz/test/lit/llvm/irreducible_loop.ll index e81e139e52..770a31740a 100644 --- a/modules/compiler/vecz/test/lit/llvm/irreducible_loop.ll +++ b/modules/compiler/vecz/test/lit/llvm/irreducible_loop.ll @@ -14,8 +14,7 @@ ; ; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception -; RUN: %pp-llvm-ver -o %t < %s --llvm-ver %LLVMVER -; RUN: veczc -k irreducible_loop -S < %s | FileCheck %t +; RUN: veczc -k irreducible_loop -S < %s | FileCheck %s ; ModuleID = 'Unknown buffer' target datalayout = "e-m:e-i64:64-f80:128-n8:16:32:64-S128" @@ -49,17 +48,10 @@ declare i64 @__mux_get_global_id(i32) ; CHECK: define spir_kernel void @__vecz_v4_irreducible_loop ; CHECK: entry: -; CHECK-LT20: br label %irr.guard.outer - -; CHECK-LT20: irr.guard.outer: ; preds = %irr.guard.pure_exit, %entry ; CHECK: br label %irr.guard -; CHECK-LT20: do.end: ; preds = %irr.guard.pure_exit -; CHECK-LT20: ret void - ; CHECK: irr.guard: ; CHECK: br i1 %{{.+}}, label %irr.guard.pure_exit, label %irr.guard ; CHECK: irr.guard.pure_exit: ; preds = %irr.guard -; CHECK-LT20: br i1 %{{.+}}, label %do.end, label %irr.guard.outer -; CHECK-GE20: ret void +; CHECK: ret void