From 4c965a09ca2b067e09db0cab059d1d2429f7341e Mon Sep 17 00:00:00 2001 From: Maksim Levental Date: Fri, 18 Oct 2024 12:41:44 -0700 Subject: [PATCH] Bump IREE to df5e5aab044ed5b6c5860b0b291c95eafe1c2522 (#848) --- .github/workflows/ci-linux.yml | 7 ++----- .github/workflows/ci-macos.yml | 9 +++++++-- .github/workflows/ci-windows.yml | 3 ++- .../Transforms/AMDAIECreateAIEWorkgroup.cpp | 4 ++-- .../Transforms/AMDAIEInsertCores.cpp | 17 ++++++++++++----- third_party/iree | 2 +- 6 files changed, 26 insertions(+), 16 deletions(-) diff --git a/.github/workflows/ci-linux.yml b/.github/workflows/ci-linux.yml index e7dcb810e..9acdcf2ec 100644 --- a/.github/workflows/ci-linux.yml +++ b/.github/workflows/ci-linux.yml @@ -59,11 +59,8 @@ jobs: - name: Python deps run: | - pip install "numpy<2" pyyaml "pybind11[global]==2.10.3" nanobind pytest - - - name: Run Pytest - run: | - pytest build_tools/ci + pip install -r third_party/iree/runtime/bindings/python/iree/runtime/build_requirements.txt + pip install pyyaml - name: Enable cache uses: actions/cache/restore@v3 diff --git a/.github/workflows/ci-macos.yml b/.github/workflows/ci-macos.yml index 79614c5b5..9d273d462 100644 --- a/.github/workflows/ci-macos.yml +++ b/.github/workflows/ci-macos.yml @@ -36,7 +36,11 @@ jobs: strategy: fail-fast: false matrix: - runs-on: [macos-12, macos-14] + runs-on: [ + macos-12, + # broken because of https://github.com/iree-org/iree/blob/0c6a151c65285987f5daabc7f76fe57a82b45ab1/compiler/plugins/target/LLVMCPU/ResolveCPUAndCPUFeatures.cpp#L58-L64 + # macos-14 + ] env: CACHE_DIR: ${{ github.workspace }}/.container-cache CACHE_KEY: ${{ matrix.runs-on }}-build-test-cpp-asserts-v1-${{ format('{0}-{1}', github.ref_name, github.run_number) }} @@ -70,7 +74,8 @@ jobs: - name: Python deps run: | - pip install "numpy<2" pyyaml "pybind11[global]==2.10.3" nanobind + pip install -r third_party/iree/runtime/bindings/python/iree/runtime/build_requirements.txt + pip install pytest - name: Enable cache uses: actions/cache/restore@v3 diff --git a/.github/workflows/ci-windows.yml b/.github/workflows/ci-windows.yml index 2fcc702e7..48777f7f5 100644 --- a/.github/workflows/ci-windows.yml +++ b/.github/workflows/ci-windows.yml @@ -77,7 +77,8 @@ jobs: - name: Python deps run: | - pip install "numpy<2" pyyaml "pybind11[global]==2.10.3" nanobind + pip install -r third_party\iree\runtime\bindings\python\iree\runtime\build_requirements.txt + pip install pyyaml - name: Enable cache uses: actions/cache/restore@v3 diff --git a/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIECreateAIEWorkgroup.cpp b/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIECreateAIEWorkgroup.cpp index 02e8fb9c0..186d16542 100644 --- a/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIECreateAIEWorkgroup.cpp +++ b/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIECreateAIEWorkgroup.cpp @@ -9,9 +9,9 @@ #include "iree-amd-aie/IR/AMDAIEDialect.h" #include "iree-amd-aie/IR/AMDAIEOps.h" #include "iree-amd-aie/Transforms/Passes.h" -#include "iree/compiler/Codegen/TransformStrategies/GPU/Common.h" #include "mlir/Dialect/Arith/Utils/Utils.h" #include "mlir/Dialect/Func/IR/FuncOps.h" +#include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/SCF/IR/SCF.h" #define DEBUG_TYPE "iree-amdaie-create-aie-workgroup" @@ -440,7 +440,7 @@ class AMDAIECreateAIEWorkgroupPass } AMDAIECreateAIEWorkgroupPass() = default; - AMDAIECreateAIEWorkgroupPass(const AMDAIECreateAIEWorkgroupPass &pass){}; + AMDAIECreateAIEWorkgroupPass(const AMDAIECreateAIEWorkgroupPass &pass) {}; void runOnOperation() override; }; diff --git a/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIEInsertCores.cpp b/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIEInsertCores.cpp index e5a797650..fd0f53325 100644 --- a/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIEInsertCores.cpp +++ b/compiler/plugins/target/AMD-AIE/iree-amd-aie/Transforms/AMDAIEInsertCores.cpp @@ -18,7 +18,8 @@ #include "iree-amd-aie/IR/AMDAIEOps.h" #include "iree-amd-aie/Transforms/AMDAIEOpUtils.h" #include "iree-amd-aie/Transforms/Passes.h" -#include "iree/compiler/Codegen/TransformStrategies/GPU/Common.h" +#include "mlir/Dialect/GPU/IR/GPUDialect.h" +#include "mlir/Dialect/Linalg/IR/Linalg.h" #include "mlir/Dialect/Linalg/IR/LinalgInterfaces.h" #include "mlir/Dialect/MemRef/IR/MemRef.h" #include "mlir/Dialect/Vector/IR/VectorOps.h" @@ -79,13 +80,19 @@ LogicalResult insertCoreOps(mlir::ModuleOp moduleOp) { auto parentOps = getInclusiveParentsOfType(forallOp); DenseMap attrMapping; getAttributeMapping(parentOps, attrMapping); - if (!attrMapping.contains(gpu::threadX(forallOp->getContext())) || - !attrMapping.contains(gpu::threadY(forallOp->getContext()))) { + mlir::gpu::GPUThreadMappingAttr threadXAttr = + mlir::gpu::GPUThreadMappingAttr::get(forallOp->getContext(), + mlir::gpu::MappingId::DimX); + mlir::gpu::GPUThreadMappingAttr threadYAttr = + mlir::gpu::GPUThreadMappingAttr::get(forallOp->getContext(), + mlir::gpu::MappingId::DimY); + if (!attrMapping.contains(threadXAttr) || + !attrMapping.contains(threadYAttr)) { forallOp.emitOpError() << "no forall with thread mapping found"; return WalkResult::interrupt(); } - Value threadX = attrMapping[gpu::threadX(forallOp->getContext())]; - Value threadY = attrMapping[gpu::threadY(forallOp->getContext())]; + Value threadX = attrMapping[threadXAttr]; + Value threadY = attrMapping[threadYAttr]; // Find input and output DMAs that need to be added to the core. SmallVector inputDmas; diff --git a/third_party/iree b/third_party/iree index 05bbcf138..df5e5aab0 160000 --- a/third_party/iree +++ b/third_party/iree @@ -1 +1 @@ -Subproject commit 05bbcf1385146d075829cd940a52bf06961614d0 +Subproject commit df5e5aab044ed5b6c5860b0b291c95eafe1c2522