From 63debf7ca51455e01b84d3dd7e6b1d7402d4b011 Mon Sep 17 00:00:00 2001 From: fschlimb Date: Tue, 10 Nov 2026 03:54:12 -0800 Subject: [PATCH 1/5] CI now covers full python matrix, using cache to not rebuilt LLVM/tileir-to-mlir SDK --- .github/workflows/ci-cpu.yml | 86 ++++++++++++++++++++++---- CMakeLists.txt | 6 ++ backends/CMakeLists.txt | 48 ++++++++++---- backends/cmake/FetchTileIRToMLIR.cmake | 4 +- backends/cmake/TileIRToMLIR.cmake | 51 +++++++++++---- backends/compiler.cmake | 32 +++++++--- backends/cpu/compiler.cmake | 7 ++- backends/test/test_compiler_build.py | 29 +++++++++ 8 files changed, 214 insertions(+), 49 deletions(-) diff --git a/.github/workflows/ci-cpu.yml b/.github/workflows/ci-cpu.yml index bcf4bbfe..ca1a580f 100644 --- a/.github/workflows/ci-cpu.yml +++ b/.github/workflows/ci-cpu.yml @@ -12,7 +12,8 @@ on: concurrency: group: ${{ github.workflow }}-${{ github.ref }} - cancel-in-progress: true + # Let main runs finish populating the native SDK cache. + cancel-in-progress: ${{ github.event_name == 'pull_request' }} permissions: contents: read @@ -46,25 +47,82 @@ jobs: - name: Check inline samples are up to date run: python test/tools/inline_samples.py --check - build: - name: Build CPU Wheel (Python 3.12) + native-sdk: + name: Native SDK runs-on: ubuntu-24.04 timeout-minutes: 360 - env: - CUTILE_BUILD_VARIANT: cpu + outputs: + key: ${{ steps.key.outputs.key }} steps: - name: Checkout repository uses: actions/checkout@v6 + - name: Compute native SDK cache key + id: key + # Bump v1 when the SDK build commands below change. + run: echo "key=cpu-native-sdk-v1-${ImageOS:?}-${RUNNER_ARCH}-${{ hashFiles('backends/cmake/**', 'backends/cpu/toolchain.cmake') }}" >> "$GITHUB_OUTPUT" + + - name: Look up native SDK cache + id: cache + uses: actions/cache/restore@v6 + with: + path: native-sdk + key: ${{ steps.key.outputs.key }} + lookup-only: true + - name: Set up Python + if: steps.cache.outputs.cache-hit != 'true' uses: actions/setup-python@v6 with: python-version: "3.12" - - name: Install native build dependencies + - name: Build native SDK + if: steps.cache.outputs.cache-hit != 'true' run: | sudo apt-get update sudo apt-get install --yes clang lld g++ + python -m pip install 'cmake>=3.24,<4' ninja + cmake -S . -B build/native-sdk -G Ninja \ + -DCUTILE_BUILD_VARIANT=cpu \ + -DCUTILE_BUILD_NATIVE_SDK=ON \ + -DCMAKE_INSTALL_PREFIX="$PWD/native-sdk" + cmake --build build/native-sdk + cmake --install build/native-sdk --component cuTileNativeSDK + + - name: Save native SDK cache + if: steps.cache.outputs.cache-hit != 'true' + uses: actions/cache/save@v6 + with: + path: native-sdk + key: ${{ steps.key.outputs.key }} + + build: + name: Build CPU Wheel (Python ${{ matrix.python-version }}) + needs: native-sdk + runs-on: ubuntu-24.04 + timeout-minutes: 60 + strategy: + fail-fast: false + matrix: + python-version: &python-versions ["3.11", "3.12", "3.13", "3.14", "3.14t", "3.15", "3.15t"] + env: + CUTILE_BUILD_VARIANT: cpu + CMAKE_ARGS: -DCUTILE_NATIVE_SDK_DIR=${{ github.workspace }}/native-sdk + steps: + - name: Checkout repository + uses: actions/checkout@v6 + + - name: Set up Python + uses: actions/setup-python@v6 + with: + python-version: ${{ matrix.python-version }} + + - name: Restore native SDK + uses: actions/cache/restore@v6 + with: + path: native-sdk + key: ${{ needs.native-sdk.outputs.key }} + fail-on-cache-miss: true - name: Set up uv uses: astral-sh/setup-uv@v7 @@ -72,7 +130,7 @@ jobs: enable-cache: true - name: Install CPU build dependencies - run: uv sync --group cpu --no-install-project + run: uv sync --group backend-build --no-install-project --python "$(command -v python)" - name: Build CPU wheel run: uv build --wheel --no-build-isolation --python .venv/bin/python @@ -80,16 +138,20 @@ jobs: - name: Upload CPU wheel uses: actions/upload-artifact@v6 with: - name: wheel-cpu-py3.12-linux-x86_64 + name: wheel-cpu-py${{ matrix.python-version }}-linux-x86_64 path: dist/*.whl if-no-files-found: error retention-days: 7 test-cpu: - name: CPU Tests (Python 3.12) + name: CPU Tests (Python ${{ matrix.python-version }}) needs: build runs-on: ubuntu-24.04 timeout-minutes: 360 + strategy: + fail-fast: false + matrix: + python-version: *python-versions steps: - name: Checkout repository uses: actions/checkout@v6 @@ -97,12 +159,12 @@ jobs: - name: Set up Python uses: actions/setup-python@v6 with: - python-version: "3.12" + python-version: ${{ matrix.python-version }} - name: Download CPU wheel uses: actions/download-artifact@v7 with: - name: wheel-cpu-py3.12-linux-x86_64 + name: wheel-cpu-py${{ matrix.python-version }}-linux-x86_64 path: dist/ - name: Install test dependencies @@ -118,6 +180,6 @@ jobs: if: ${{ !cancelled() }} uses: actions/upload-artifact@v6 with: - name: cpu-test-results + name: cpu-test-results-py${{ matrix.python-version }} path: test-results.xml retention-days: 7 \ No newline at end of file diff --git a/CMakeLists.txt b/CMakeLists.txt index 1df6311e..bc3465fd 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -6,6 +6,12 @@ cmake_minimum_required(VERSION 3.18) project(cuda-tile-python) enable_testing() +option(CUTILE_BUILD_NATIVE_SDK "Build and install the Python-independent native SDK" OFF) +if(CUTILE_BUILD_NATIVE_SDK) + add_subdirectory(backends) + return() +endif() + list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_LIST_DIR}/cmake") if (MSVC) diff --git a/backends/CMakeLists.txt b/backends/CMakeLists.txt index 316d2b80..fb9618ce 100644 --- a/backends/CMakeLists.txt +++ b/backends/CMakeLists.txt @@ -26,6 +26,14 @@ endif() # dependencies, such as dlpack. set(CMAKE_INSTALL_DEFAULT_COMPONENT_NAME cuTileBackend) +set(CUTILE_NATIVE_SDK_DIR "" CACHE PATH "Installed Python-independent native SDK") +if(CUTILE_BUILD_NATIVE_SDK) + include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/FetchTileIRToMLIR.cmake") + fetch_tileir_to_mlir() + include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/TileIRToMLIR.cmake") + return() +endif() + find_package(Python REQUIRED COMPONENTS Interpreter Development.Module) execute_process( COMMAND "${Python_EXECUTABLE}" -m nanobind --cmake_dir @@ -37,16 +45,34 @@ find_package(nanobind CONFIG REQUIRED) # Package directory of the backends, relative to the install prefix. set(CUTILE_BACKEND_INSTALL_DIR cuda/tile/_backends) -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/FetchTileIRToMLIR.cmake") -fetch_tileir_to_mlir() -include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/TileIRToMLIR.cmake") - -install(FILES "${CUTILE_COMPILER_PATH}" DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}") -install(FILES "${CUTILE_LLVM_LICENSE_PATH}" - DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/llvm") -install(FILES "${CUTILE_TILEIR_TO_MLIR_DIR}/LICENSE" - DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/tileir-to-mlir") -install(FILES "${CUTILE_CUDA_TILE_SOURCE_DIR}/LICENSE.txt" - DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/cuda-tile") +if(CUTILE_NATIVE_SDK_DIR) + list(PREPEND CMAKE_PREFIX_PATH "${CUTILE_NATIVE_SDK_DIR}") + find_package(TileIRToMLIR REQUIRED CONFIG + PATHS "${CUTILE_NATIVE_SDK_DIR}" NO_DEFAULT_PATH) + find_package(LLD REQUIRED CONFIG) + add_library(TileIRToMLIR ALIAS TileIRToMLIR::TileIRToMLIR) + if(NOT TARGET CudaTileBytecodeReader) + add_library(CudaTileBytecodeReader ALIAS TileIRToMLIR::CudaTileBytecodeReader) + endif() + include("${CMAKE_CURRENT_SOURCE_DIR}/compiler.cmake") + install(FILES "${CUTILE_NATIVE_SDK_DIR}/share/cutile/llvm/LICENSE.TXT" + DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/llvm") + install(FILES "${CUTILE_NATIVE_SDK_DIR}/share/TileIRToMLIR/LICENSE" + DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/tileir-to-mlir") + install(FILES "${CUTILE_NATIVE_SDK_DIR}/share/TileIRToMLIR/cuda-tile/LICENSE.txt" + DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/cuda-tile") +else() + include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/FetchTileIRToMLIR.cmake") + fetch_tileir_to_mlir() + include("${CMAKE_CURRENT_SOURCE_DIR}/cmake/TileIRToMLIR.cmake") + + install(FILES "${CUTILE_COMPILER_PATH}" DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}") + install(FILES "${CUTILE_LLVM_LICENSE_PATH}" + DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/llvm") + install(FILES "${CUTILE_TILEIR_TO_MLIR_DIR}/LICENSE" + DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/tileir-to-mlir") + install(FILES "${CUTILE_CUDA_TILE_SOURCE_DIR}/LICENSE.txt" + DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}/_licenses/cuda-tile") +endif() add_subdirectory("${CUTILE_BUILD_VARIANT}") diff --git a/backends/cmake/FetchTileIRToMLIR.cmake b/backends/cmake/FetchTileIRToMLIR.cmake index 9ed4b258..477ab752 100644 --- a/backends/cmake/FetchTileIRToMLIR.cmake +++ b/backends/cmake/FetchTileIRToMLIR.cmake @@ -4,10 +4,10 @@ set(TILEIR_TO_MLIR_REPOSITORY "https://github.com/libxsmm/tileir-to-mlir" CACHE STRING "TileIRToMLIR repo URL") -set(TILEIR_TO_MLIR_REVISION "ba6c9e50274522089b7ad99b3e8a571480481330" +set(TILEIR_TO_MLIR_REVISION "24eadac04f384e7bd00fa9533974537474a05082" CACHE STRING "TileIRToMLIR revision") set(TILEIR_TO_MLIR_REVISION_SHA256 - "f39644ec0eb4d98db8de6b8bcf51f82d67ea80081080e6c397695a20369eb0d4" + "76755bf0f7b1ff8eaeab39297bc9282e52679b146af9367113999a76e77fa502" CACHE STRING "Expected SHA256 of the pinned TileIRToMLIR archive") set(CUDA_TILE_REPOSITORY "https://github.com/NVIDIA/cuda-tile" diff --git a/backends/cmake/TileIRToMLIR.cmake b/backends/cmake/TileIRToMLIR.cmake index a40e62f7..b86f4d36 100644 --- a/backends/cmake/TileIRToMLIR.cmake +++ b/backends/cmake/TileIRToMLIR.cmake @@ -5,12 +5,14 @@ set(CUTILE_TOOLCHAIN_DIR "" CACHE PATH "Existing native compiler installation for CUTILE_BUILD_VARIANT; skips building LLVM") -execute_process( - COMMAND "${Python_EXECUTABLE}" -c - "import sysconfig; print(sysconfig.get_config_var('EXT_SUFFIX'))" - OUTPUT_VARIABLE _compiler_suffix OUTPUT_STRIP_TRAILING_WHITESPACE - COMMAND_ERROR_IS_FATAL ANY) -set(_compiler_name "_compiler${_compiler_suffix}") +if(NOT CUTILE_BUILD_NATIVE_SDK) + execute_process( + COMMAND "${Python_EXECUTABLE}" -c + "import sysconfig; print(sysconfig.get_config_var('EXT_SUFFIX'))" + OUTPUT_VARIABLE _compiler_suffix OUTPUT_STRIP_TRAILING_WHITESPACE + COMMAND_ERROR_IS_FATAL ANY) + set(_compiler_name "_compiler${_compiler_suffix}") +endif() set(_tileir_to_mlir_dir "${CUTILE_TILEIR_TO_MLIR_DIR}") if(NOT EXISTS "${_tileir_to_mlir_dir}/CMakeLists.txt") @@ -18,7 +20,7 @@ if(NOT EXISTS "${_tileir_to_mlir_dir}/CMakeLists.txt") "TileIRToMLIR sources are missing: ${_tileir_to_mlir_dir}") endif() -if(CUTILE_TOOLCHAIN_DIR) +if(CUTILE_TOOLCHAIN_DIR AND NOT CUTILE_BUILD_NATIVE_SDK) set(CUTILE_COMPILER_PATH "${CUTILE_TOOLCHAIN_DIR}/lib/${_compiler_name}") set(CUTILE_LLVM_LICENSE_PATH "${CUTILE_TOOLCHAIN_DIR}/share/cutile/llvm/LICENSE.TXT") foreach(_path IN ITEMS CUTILE_COMPILER_PATH CUTILE_LLVM_LICENSE_PATH) @@ -77,6 +79,15 @@ set(_cutile_checkout_command include("${CMAKE_CURRENT_SOURCE_DIR}/${CUTILE_BUILD_VARIANT}/toolchain.cmake") set(_cutile_llvm_build_targets _compiler) set(_cutile_llvm_byproducts "${CUTILE_COMPILER_PATH}") +set(_cutile_external_projects "TileIRToMLIR|cuTileCompiler") +if(CUTILE_BUILD_NATIVE_SDK) + set(_cutile_external_projects TileIRToMLIR) + set(_cutile_llvm_build_targets install-distribution) + set(_cutile_llvm_byproducts + "${CMAKE_INSTALL_PREFIX}/lib/cmake/TileIRToMLIR/TileIRToMLIRConfig.cmake") + install(FILES "${CUTILE_LLVM_LICENSE_PATH}" + DESTINATION share/cutile/llvm COMPONENT cuTileNativeSDK) +endif() set(_cutile_llvm_cmake_args "-DCMAKE_BUILD_TYPE:STRING=${CUTILE_LLVM_BUILD_TYPE}" @@ -84,7 +95,7 @@ set(_cutile_llvm_cmake_args "-DCMAKE_CXX_COMPILER:FILEPATH=${CUTILE_CLANGXX_EXECUTABLE}" "-DCMAKE_LINKER:FILEPATH=${CUTILE_LLD_EXECUTABLE}" "-DLLVM_ENABLE_PROJECTS:STRING=${_cutile_llvm_projects}" - "-DLLVM_EXTERNAL_PROJECTS:STRING=TileIRToMLIR|cuTileCompiler" + "-DLLVM_EXTERNAL_PROJECTS:STRING=${_cutile_external_projects}" "-DLLVM_EXTERNAL_TILEIRTOMLIR_SOURCE_DIR:PATH=${_tileir_to_mlir_dir}" # Explicit, so that a shared LLVM tree never keeps another checkout's path. "-DTILEIR_TO_MLIR_CUDA_TILE_SOURCE_DIR:PATH=${CUTILE_CUDA_TILE_SOURCE_DIR}" @@ -114,10 +125,20 @@ set(_cutile_llvm_cmake_args -DMLIR_INCLUDE_TESTS:BOOL=OFF -DMLIR_ENABLE_EXECUTION_ENGINE:BOOL=ON -DMLIR_ENABLE_LEVELZERO_RUNNER:BOOL=OFF - -DMLIR_ENABLE_BINDINGS_PYTHON:BOOL=OFF - "-DPython3_EXECUTABLE:FILEPATH=${Python_EXECUTABLE}" - "-DPython_EXECUTABLE:FILEPATH=${Python_EXECUTABLE}" - "-Dnanobind_DIR:PATH=${nanobind_DIR}") + -DMLIR_ENABLE_BINDINGS_PYTHON:BOOL=OFF) +if(CUTILE_BUILD_NATIVE_SDK) + list(APPEND _cutile_llvm_cmake_args + -DLLVM_ENABLE_LIBEDIT:BOOL=OFF + "-DCMAKE_INSTALL_PREFIX:PATH=${CMAKE_INSTALL_PREFIX}" + -DCMAKE_INSTALL_LIBDIR:STRING=lib + "-DLLVM_DISTRIBUTION_COMPONENTS:STRING=llvm-libraries|mlir-libraries|lld-libraries|llvm-headers|mlir-headers|lld-headers|cmake-exports|mlir-cmake-exports|lld-cmake-exports|TileIRToMLIR") +else() + list(APPEND _cutile_llvm_cmake_args + -DLLVM_DISTRIBUTION_COMPONENTS:STRING= + "-DPython3_EXECUTABLE:FILEPATH=${Python_EXECUTABLE}" + "-DPython_EXECUTABLE:FILEPATH=${Python_EXECUTABLE}" + "-Dnanobind_DIR:PATH=${nanobind_DIR}") +endif() list(APPEND _cutile_llvm_cmake_args ${_cutile_backend_cmake_args}) set(_cutile_llvm_build_script "${CMAKE_BINARY_DIR}/tileir-toolchain/BuildLLVM.cmake") @@ -149,7 +170,11 @@ else() message(STATUS "LLVM revision: ${_tileir_required_llvm_revision}") endif() message(STATUS "LLVM source: ${CUTILE_LLVM_SOURCE_DIR}") -message(STATUS "Native cuTile compiler: ${CUTILE_COMPILER_PATH}") +if(CUTILE_BUILD_NATIVE_SDK) + message(STATUS "Native SDK: ${CMAKE_INSTALL_PREFIX}") +else() + message(STATUS "Native cuTile compiler: ${CUTILE_COMPILER_PATH}") +endif() unset(_cutile_checkout_command) unset(_cutile_checkout_script) diff --git a/backends/compiler.cmake b/backends/compiler.cmake index 53c3be1d..afe8dd91 100644 --- a/backends/compiler.cmake +++ b/backends/compiler.cmake @@ -11,9 +11,15 @@ find_package(nanobind CONFIG REQUIRED) nanobind_add_module(_compiler FREE_THREADED mlir/compiler.cpp) target_include_directories(_compiler PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}") -target_include_directories(_compiler SYSTEM PRIVATE - "${LLVM_MAIN_SRC_DIR}/../mlir/include" - "${LLVM_BINARY_DIR}/tools/mlir/include") +if(CUTILE_NATIVE_SDK_DIR) + target_include_directories(_compiler SYSTEM PRIVATE ${LLVM_INCLUDE_DIRS} ${MLIR_INCLUDE_DIRS}) + separate_arguments(_llvm_definitions NATIVE_COMMAND "${LLVM_DEFINITIONS}") + target_compile_definitions(_compiler PRIVATE ${_llvm_definitions}) +else() + target_include_directories(_compiler SYSTEM PRIVATE + "${LLVM_MAIN_SRC_DIR}/../mlir/include" + "${LLVM_BINARY_DIR}/tools/mlir/include") +endif() target_link_libraries(_compiler PRIVATE TileIRToMLIR CudaTileBytecodeReader MLIRParser MLIRTransforms @@ -29,9 +35,13 @@ target_link_libraries(_compiler PRIVATE MLIRVectorTransforms) target_sources(_compiler PRIVATE mlir/LowerVectorScan.cpp) target_link_options(_compiler PRIVATE "LINKER:--exclude-libs,ALL") -set_target_properties(_compiler PROPERTIES - LIBRARY_OUTPUT_DIRECTORY "${LLVM_BINARY_DIR}/lib" - BUILD_WITH_INSTALL_RPATH TRUE) +set_target_properties(_compiler PROPERTIES BUILD_WITH_INSTALL_RPATH TRUE) +if(CUTILE_NATIVE_SDK_DIR) + set_target_properties(_compiler PROPERTIES + LIBRARY_OUTPUT_DIRECTORY "${CMAKE_CURRENT_BINARY_DIR}/lib") +else() + set_target_properties(_compiler PROPERTIES LIBRARY_OUTPUT_DIRECTORY "${LLVM_BINARY_DIR}/lib") +endif() include("${CMAKE_CURRENT_SOURCE_DIR}/${CUTILE_BUILD_VARIANT}/compiler.cmake") @@ -41,6 +51,10 @@ if(NOT LLVM_ENABLE_RTTI) PROPERTIES COMPILE_OPTIONS "-fno-rtti") endif() -install(TARGETS _compiler LIBRARY DESTINATION lib COMPONENT cuTileCompiler) -install(FILES "${LLVM_MAIN_SRC_DIR}/LICENSE.TXT" - DESTINATION share/cutile/llvm COMPONENT cuTileCompiler) +if(CUTILE_NATIVE_SDK_DIR) + install(TARGETS _compiler LIBRARY DESTINATION "${CUTILE_BACKEND_INSTALL_DIR}") +else() + install(TARGETS _compiler LIBRARY DESTINATION lib COMPONENT cuTileCompiler) + install(FILES "${LLVM_MAIN_SRC_DIR}/LICENSE.TXT" + DESTINATION share/cutile/llvm COMPONENT cuTileCompiler) +endif() diff --git a/backends/cpu/compiler.cmake b/backends/cpu/compiler.cmake index 7756fde2..28b255e2 100644 --- a/backends/cpu/compiler.cmake +++ b/backends/cpu/compiler.cmake @@ -3,8 +3,11 @@ # SPDX-License-Identifier: Apache-2.0 target_sources(_compiler PRIVATE cpu/compiler_cpu.cpp cpu/ApproximateMath.cpp) -target_include_directories(_compiler SYSTEM PRIVATE - "${LLVM_MAIN_SRC_DIR}/../lld/include") +if(CUTILE_NATIVE_SDK_DIR) + target_include_directories(_compiler SYSTEM PRIVATE ${LLD_INCLUDE_DIRS}) +else() + target_include_directories(_compiler SYSTEM PRIVATE "${LLVM_MAIN_SRC_DIR}/../lld/include") +endif() target_link_libraries(_compiler PRIVATE MLIRMemRefTransforms MLIRMathTransforms MLIRAffineToStandard MLIRVectorToSCF diff --git a/backends/test/test_compiler_build.py b/backends/test/test_compiler_build.py index 7a7a2bd9..7c9d2119 100644 --- a/backends/test/test_compiler_build.py +++ b/backends/test/test_compiler_build.py @@ -187,3 +187,32 @@ def test_cpu_toolchain_settings(tmp_path): """) subprocess.run([cmake, "-P", str(script)], check=True, capture_output=True, text=True, timeout=60) + + +def test_native_sdk_configure_without_python_or_cuda(compiler_sources, tmp_path): + if any(shutil.which(tool) is None for tool in ("clang", "clang++", "ld.lld")): + pytest.skip("native SDK configuration requires Clang and LLD") + run, binary, _ = compiler_sources + run() + tileir, cuda_tile = (binary / "sources.txt").read_text().splitlines() + project = Path(__file__).resolve().parents[2] + output = tmp_path / "sdk-build" + prefix = tmp_path / "sdk-install" + subprocess.run([ + shutil.which("cmake"), "-S", str(project), "-B", str(output), + "-DCUTILE_BUILD_VARIANT=cpu", "-DCUTILE_BUILD_NATIVE_SDK=ON", + "-DCMAKE_DISABLE_FIND_PACKAGE_Python=ON", + "-DCMAKE_DISABLE_FIND_PACKAGE_CUDAToolkit=ON", + f"-DCMAKE_INSTALL_PREFIX={prefix}", + f"-DCUTILE_LLVM_SOURCE_DIR={tmp_path / 'llvm-source'}", + f"-DCUTILE_LLVM_BINARY_DIR={tmp_path / 'llvm-build'}", + f"-DFETCHCONTENT_SOURCE_DIR_TILEIR_TO_MLIR={tileir}", + f"-DFETCHCONTENT_SOURCE_DIR_CUDA_TILE={cuda_tile}", + ], check=True, capture_output=True, text=True, timeout=60) + script = (output / "tileir-toolchain" / "BuildLLVM.cmake").read_text() + assert "set(_targets [==[install-distribution]==])" in script + assert "-DLLVM_EXTERNAL_PROJECTS:STRING=TileIRToMLIR;" in script + assert f"-DCMAKE_INSTALL_PREFIX:PATH={prefix}" in script + assert "-DPython_EXECUTABLE" not in script + assert "-Dnanobind_DIR" not in script + assert "-DLLVM_ENABLE_LIBEDIT:BOOL=OFF" in script From fee8d10e30ebc8faa7aacddd56e5de558302b887 Mon Sep 17 00:00:00 2001 From: fschlimb Date: Tue, 10 Nov 2026 06:22:47 -0800 Subject: [PATCH 2/5] fix for ci-cpu --- .github/workflows/ci-cpu.yml | 15 +++++++++++++++ backends/cmake/FetchTileIRToMLIR.cmake | 4 ++-- 2 files changed, 17 insertions(+), 2 deletions(-) diff --git a/.github/workflows/ci-cpu.yml b/.github/workflows/ci-cpu.yml index ca1a580f..0211677e 100644 --- a/.github/workflows/ci-cpu.yml +++ b/.github/workflows/ci-cpu.yml @@ -96,6 +96,16 @@ jobs: path: native-sdk key: ${{ steps.key.outputs.key }} + - name: Verify native SDK cache + if: steps.cache.outputs.cache-hit != 'true' + # Save failures are only warnings; the wheel jobs need the entry. + uses: actions/cache/restore@v6 + with: + path: native-sdk + key: ${{ steps.key.outputs.key }} + lookup-only: true + fail-on-cache-miss: true + build: name: Build CPU Wheel (Python ${{ matrix.python-version }}) needs: native-sdk @@ -173,6 +183,11 @@ jobs: python -m pip install dist/*.whl numpy pytest pytest-benchmark pytest-env python -m pip install setuptools==84.0.0 'cmake>=3.24,<4' ninja + - name: Check native extensions keep the GIL disabled + if: endsWith(matrix.python-version, 't') + # CPython only warns when an extension re-enables the GIL. + run: python -c "import sys, cuda.tile._cext, cuda.tile._backends._compiler, cuda.tile._backends.cpu._cpu; assert not sys._is_gil_enabled(), 'GIL re-enabled'" + - name: Run CPU backend tests run: python -m pytest --durations=10 --junitxml=test-results.xml -m "not benchmark" backends/cpu backends/test diff --git a/backends/cmake/FetchTileIRToMLIR.cmake b/backends/cmake/FetchTileIRToMLIR.cmake index 477ab752..4967075c 100644 --- a/backends/cmake/FetchTileIRToMLIR.cmake +++ b/backends/cmake/FetchTileIRToMLIR.cmake @@ -4,10 +4,10 @@ set(TILEIR_TO_MLIR_REPOSITORY "https://github.com/libxsmm/tileir-to-mlir" CACHE STRING "TileIRToMLIR repo URL") -set(TILEIR_TO_MLIR_REVISION "24eadac04f384e7bd00fa9533974537474a05082" +set(TILEIR_TO_MLIR_REVISION "7eeb1655cb4af986cc5aa53ca35ff87e7770d005" CACHE STRING "TileIRToMLIR revision") set(TILEIR_TO_MLIR_REVISION_SHA256 - "76755bf0f7b1ff8eaeab39297bc9282e52679b146af9367113999a76e77fa502" + "c3946e02afeb443d4961dd81739f7d3fd3b94383871490df6f2cb2153b513c00" CACHE STRING "Expected SHA256 of the pinned TileIRToMLIR archive") set(CUDA_TILE_REPOSITORY "https://github.com/NVIDIA/cuda-tile" From 949190b8aaedeaab2de11e97fc8a9bf70c2ab3bd Mon Sep 17 00:00:00 2001 From: fschlimb Date: Tue, 10 Nov 2026 08:19:51 -0800 Subject: [PATCH 3/5] fix include order for python headers --- cext/enums.h | 4 ++-- cext/xla_ffi.cpp | 2 +- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/cext/enums.h b/cext/enums.h index d14b00a1..31680ec5 100644 --- a/cext/enums.h +++ b/cext/enums.h @@ -4,13 +4,13 @@ #pragma once +#include + #include #include #include #include -#include - #define ENUM_ENTRY_AUTO(name, ...) name, #define ENUM_ENTRY_WITH_ID(name, id, ...) name = id, diff --git a/cext/xla_ffi.cpp b/cext/xla_ffi.cpp index 4fe4b440..6e77d9c4 100644 --- a/cext/xla_ffi.cpp +++ b/cext/xla_ffi.cpp @@ -20,9 +20,9 @@ const XLA_FFI_TypeInfo* cutile_call_state_type_info() { #else // !_WIN32 +#include "py.h" #include "xla/ffi/api/c_api.h" #include "xla_ffi.h" -#include "py.h" #include "cuda_loader.h" #include "hash_map.h" #include "memory.h" From 0343be9220c58da3f5505bd1c89a7f44752c2583 Mon Sep 17 00:00:00 2001 From: fschlimb Date: Tue, 10 Nov 2026 08:56:12 -0800 Subject: [PATCH 4/5] fix include order for python headers --- backends/cpu/compiler_cpu.cpp | 4 +++- backends/mlir/compiler.h | 6 ++++-- 2 files changed, 7 insertions(+), 3 deletions(-) diff --git a/backends/cpu/compiler_cpu.cpp b/backends/cpu/compiler_cpu.cpp index fc943caf..1ce18cb5 100644 --- a/backends/cpu/compiler_cpu.cpp +++ b/backends/cpu/compiler_cpu.cpp @@ -3,6 +3,9 @@ // // SPDX-License-Identifier: Apache-2.0 +// Includes Python.h, which must precede all system headers. +#include "mlir/compiler.h" + #include "ApproximateMath.h" #include "amx.h" #include "lld/Common/Driver.h" @@ -17,7 +20,6 @@ #include "mlir/Pass/PassRegistry.h" #include "mlir/Target/LLVMIR/Export.h" #include "mlir/Transforms/Passes.h" -#include "mlir/compiler.h" #include "llvm/IR/LegacyPassManager.h" #include "llvm/IR/Module.h" #include "llvm/IR/Verifier.h" diff --git a/backends/mlir/compiler.h b/backends/mlir/compiler.h index 1bafa3fc..eead87ec 100644 --- a/backends/mlir/compiler.h +++ b/backends/mlir/compiler.h @@ -5,6 +5,10 @@ #pragma once +// nanobind includes Python.h, which must precede all system headers. +#include +#include + #include "mlir/IR/BuiltinOps.h" #include "mlir/IR/Diagnostics.h" #include "mlir/IR/MLIRContext.h" @@ -12,8 +16,6 @@ #include "llvm/ADT/STLFunctionalExtras.h" #include "llvm/ADT/StringRef.h" #include "llvm/Support/raw_ostream.h" -#include -#include #include namespace cutile { From 115d00c997e88fb15b4bb9c8733740c10790be7f Mon Sep 17 00:00:00 2001 From: fschlimb Date: Wed, 11 Nov 2026 02:43:20 -0800 Subject: [PATCH 5/5] Backport LLVM fix for bf16 broadcasts on AVX2 Assisted by Copilot --- backends/cmake/CheckoutLLVM.cmake | 49 ++++++++++++++++++++++-- backends/cmake/TileIRToMLIR.cmake | 5 ++- backends/cmake/llvm.patch | 56 ++++++++++++++++++++++++++++ backends/test/test_compiler_build.py | 47 +++++++++++++++++++++++ 4 files changed, 152 insertions(+), 5 deletions(-) create mode 100644 backends/cmake/llvm.patch diff --git a/backends/cmake/CheckoutLLVM.cmake b/backends/cmake/CheckoutLLVM.cmake index d9dd87d5..dc3c5145 100644 --- a/backends/cmake/CheckoutLLVM.cmake +++ b/backends/cmake/CheckoutLLVM.cmake @@ -56,16 +56,57 @@ if(DEFINED REVISION AND NOT REVISION STREQUAL "") set(_selected_revision "${REVISION}") endif() +set(_patch "") +if(DEFINED PATCH AND NOT PATCH STREQUAL "") + file(READ "${PATCH}" _patch) +endif() +execute_process( + COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" rev-parse --absolute-git-dir + OUTPUT_VARIABLE _git_dir + OUTPUT_STRIP_TRAILING_WHITESPACE + COMMAND_ERROR_IS_FATAL ANY) +# Copy of the patch applied to the working tree, needed to revert it. +set(_applied_patch "${_git_dir}/cutile-llvm.patch") +set(_applied "") +if(EXISTS "${_applied_patch}") + file(READ "${_applied_patch}" _applied) +endif() + execute_process( COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" rev-parse HEAD RESULT_VARIABLE _head_result OUTPUT_VARIABLE _head_revision OUTPUT_STRIP_TRAILING_WHITESPACE ERROR_QUIET) -if(_head_result EQUAL 0 AND _head_revision STREQUAL _selected_revision) +if(_head_result EQUAL 0 AND _head_revision STREQUAL _selected_revision + AND _applied STREQUAL _patch) return() endif() -execute_process( - COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" checkout --detach "${_selected_revision}" - COMMAND_ERROR_IS_FATAL ANY) +if(EXISTS "${_applied_patch}") + execute_process( + COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" apply --reverse "${_applied_patch}" + COMMAND_ERROR_IS_FATAL ANY) + file(REMOVE "${_applied_patch}") +endif() + +if(NOT _head_revision STREQUAL _selected_revision) + execute_process( + COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" checkout --detach "${_selected_revision}" + COMMAND_ERROR_IS_FATAL ANY) +endif() + +if(NOT _patch STREQUAL "") + # Newer LLVM revisions may already contain the patch. + execute_process( + COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" apply --reverse --check "${PATCH}" + RESULT_VARIABLE _patch_present + OUTPUT_QUIET + ERROR_QUIET) + if(NOT _patch_present EQUAL 0) + execute_process( + COMMAND "${GIT_EXECUTABLE}" -C "${SOURCE_DIR}" apply "${PATCH}" + COMMAND_ERROR_IS_FATAL ANY) + file(WRITE "${_applied_patch}" "${_patch}") + endif() +endif() diff --git a/backends/cmake/TileIRToMLIR.cmake b/backends/cmake/TileIRToMLIR.cmake index b86f4d36..101cebce 100644 --- a/backends/cmake/TileIRToMLIR.cmake +++ b/backends/cmake/TileIRToMLIR.cmake @@ -68,12 +68,14 @@ file(CONFIGURE OUTPUT "${_required_llvm_revisions_file}" set(CUTILE_COMPILER_PATH "${CUTILE_LLVM_BINARY_DIR}/lib/${_compiler_name}") set(_cutile_checkout_script "${CMAKE_CURRENT_LIST_DIR}/CheckoutLLVM.cmake") +set(_cutile_llvm_patch "${CMAKE_CURRENT_LIST_DIR}/llvm.patch") set(_cutile_checkout_command "${CMAKE_COMMAND}" "-DGIT_EXECUTABLE=${GIT_EXECUTABLE}" "-DSOURCE_DIR=${CUTILE_LLVM_SOURCE_DIR}" "-DTILEIR_REQUIRED_LLVM_REVISION=${_tileir_required_llvm_revision}" "-DREVISION=${_cutile_llvm_override}" + "-DPATCH=${_cutile_llvm_patch}" -P "${_cutile_checkout_script}") include("${CMAKE_CURRENT_SOURCE_DIR}/${CUTILE_BUILD_VARIANT}/toolchain.cmake") @@ -162,7 +164,7 @@ ExternalProject_Add(cutile_llvm USES_TERMINAL_BUILD TRUE) ExternalProject_Add_StepDependencies(cutile_llvm download - "${_required_llvm_revisions_file}") + "${_required_llvm_revisions_file}" "${_cutile_llvm_patch}") if(_cutile_llvm_override) message(STATUS "LLVM revision override: ${_cutile_llvm_override}") @@ -178,6 +180,7 @@ endif() unset(_cutile_checkout_command) unset(_cutile_checkout_script) +unset(_cutile_llvm_patch) unset(_required_llvm_revisions_file) unset(_cutile_llvm_override) unset(_tileir_required_llvm_revision) diff --git a/backends/cmake/llvm.patch b/backends/cmake/llvm.patch new file mode 100644 index 00000000..0292290b --- /dev/null +++ b/backends/cmake/llvm.patch @@ -0,0 +1,56 @@ +SPDX-FileCopyrightText: Copyright (c) <2026> LLVM Project contributors +SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +Backport of https://github.com/llvm/llvm-project/pull/240805 (5f93c6a7) without +tests: AVX2 hosts without AVX512BW+VL failed to select bf16 broadcasts. +Remove when the pinned LLVM revision includes it. + +diff --git a/llvm/lib/Target/X86/X86ISelLowering.cpp b/llvm/lib/Target/X86/X86ISelLowering.cpp +index 238250a8b67a..d113e8c416f6 100644 +--- a/llvm/lib/Target/X86/X86ISelLowering.cpp ++++ b/llvm/lib/Target/X86/X86ISelLowering.cpp +@@ -43633,7 +43633,10 @@ static SDValue combineTargetShuffle(SDValue N, const SDLoc &DL, + BCVT.getScalarType().getTypeForEVT(*DAG.getContext()))) { + EVT NewVT = EVT::getVectorVT(*DAG.getContext(), BCVT.getScalarType(), + VT.getVectorNumElements()); +- return DAG.getBitcast(VT, DAG.getNode(X86ISD::VBROADCAST, DL, NewVT, BC)); ++ // NewVT can be illegal even if BCVT is legal (e.g. v16bf16 vs. v8bf16). ++ if (TLI.isTypeLegal(NewVT)) ++ return DAG.getBitcast(VT, ++ DAG.getNode(X86ISD::VBROADCAST, DL, NewVT, BC)); + } + + // vbroadcast(bitcast(vbroadcast(src))) -> bitcast(vbroadcast(src)) +@@ -43648,7 +43651,9 @@ static SDValue combineTargetShuffle(SDValue N, const SDLoc &DL, + MVT NewVT = + MVT::getVectorVT(BCVT.getSimpleVT().getScalarType(), + VT.getSizeInBits() / BCVT.getScalarSizeInBits()); +- return DAG.getBitcast(VT, DAG.getNode(X86ISD::VBROADCAST, DL, NewVT, BC)); ++ if (TLI.isTypeLegal(NewVT)) ++ return DAG.getBitcast(VT, ++ DAG.getNode(X86ISD::VBROADCAST, DL, NewVT, BC)); + } + + // Reduce broadcast source vector to lowest 128-bits. +diff --git a/llvm/lib/Target/X86/X86InstrSSE.td b/llvm/lib/Target/X86/X86InstrSSE.td +index 95426695f7b2..00b393376a40 100644 +--- a/llvm/lib/Target/X86/X86InstrSSE.td ++++ b/llvm/lib/Target/X86/X86InstrSSE.td +@@ -7852,6 +7852,17 @@ let Predicates = [HasAVX2, NoVLX_Or_NoBWI] in { + (VPBROADCASTWrr (COPY_TO_REGCLASS FR16:$src, VR128))>; + def : Pat<(v16f16 (X86VBroadcast (f16 FR16:$src))), + (VPBROADCASTWYrr (COPY_TO_REGCLASS FR16:$src, VR128))>; ++ ++ // No scalar bf16 broadcast patterns: scalar bf16 is not a legal type. ++ def : Pat<(v8bf16 (X86VBroadcastld16 addr:$src)), ++ (VPBROADCASTWrm addr:$src)>; ++ def : Pat<(v16bf16 (X86VBroadcastld16 addr:$src)), ++ (VPBROADCASTWYrm addr:$src)>; ++ ++ def : Pat<(v8bf16 (X86VBroadcast (v8bf16 VR128:$src))), ++ (VPBROADCASTWrr VR128:$src)>; ++ def : Pat<(v16bf16 (X86VBroadcast (v8bf16 VR128:$src))), ++ (VPBROADCASTWYrr VR128:$src)>; + } + let Predicates = [HasAVX2, NoVLX] in { + def : Pat<(v4i32 (X86VBroadcast GR32:$src)), diff --git a/backends/test/test_compiler_build.py b/backends/test/test_compiler_build.py index 7c9d2119..5e75817a 100644 --- a/backends/test/test_compiler_build.py +++ b/backends/test/test_compiler_build.py @@ -170,6 +170,53 @@ def test_missing_cache_triggers_reconfiguration(llvm_build): "cpu:mlir;lld", "cpu:mlir;lld"] +def test_llvm_checkout_applies_and_reverts_patch(tmp_path): + cmake, git_executable = shutil.which("cmake"), shutil.which("git") + if cmake is None or git_executable is None: + pytest.skip("LLVM checkout tests require CMake and Git") + source = tmp_path / "llvm" + source.mkdir() + target = source / "x86.td" + + def git(*args): + return subprocess.run([git_executable, "-C", str(source), *args], check=True, + capture_output=True, text=True).stdout.strip() + + def commit(text): + target.write_text(text) + git("add", target.name) + git("-c", "user.name=test", "-c", "user.email=test@example.com", + "-c", "commit.gpgsign=false", "commit", "-q", "-m", "update") + return git("rev-parse", "HEAD") + + git("init", "-q") + pinned = commit("a\nb\nc\n") + newer = commit("a\nb\nc\nd\n") + fixed = commit("a\nB\nc\nd\n") + patch = tmp_path / "llvm.patch" + patch.write_text("Backport.\n\n" + "diff --git a/x86.td b/x86.td\n--- a/x86.td\n+++ b/x86.td\n" + "@@ -1,3 +1,3 @@\n a\n-b\n+B\n c\n") + script = Path(__file__).resolve().parents[1] / "cmake" / "CheckoutLLVM.cmake" + + def checkout(revision="", patch_path=patch): + subprocess.run([cmake, f"-DGIT_EXECUTABLE={git_executable}", f"-DSOURCE_DIR={source}", + f"-DTILEIR_REQUIRED_LLVM_REVISION={pinned}", f"-DREVISION={revision}", + f"-DPATCH={patch_path}", "-P", str(script)], + check=True, capture_output=True, text=True, timeout=60) + return git("rev-parse", "HEAD"), target.read_text() + + assert checkout() == (pinned, "a\nB\nc\n") + assert checkout() == (pinned, "a\nB\nc\n") + # The patch must be reverted before checking out a revision that changes its files. + assert checkout(newer) == (newer, "a\nB\nc\nd\n") + assert checkout(fixed) == (fixed, "a\nB\nc\nd\n") + assert git("status", "--porcelain") == "" + assert checkout() == (pinned, "a\nB\nc\n") + assert checkout(patch_path="") == (pinned, "a\nb\nc\n") + assert git("status", "--porcelain") == "" + + def test_cpu_toolchain_settings(tmp_path): cmake = shutil.which("cmake") if cmake is None: