[mlir][acc] Add loop tiling utilities for OpenACC #171490
Merged
Conversation
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Add utilities in OpenACCUtilsTiling.h/.cpp to support tiling transformations on acc.loop operations: - uncollapseLoops: Expand collapsed loops with multiple IVs into nested loop structures when tile count exceeds collapse count - tileACCLoops: Transform loop nests into tile and element loops based on provided tile sizes, with automatic resolution of unknown tile sizes (tile(*) represented as -1) These utilities prepare for the ACCLoopTiling pass which handles the OpenACC loop tile directive.
Member
|
@llvm/pr-subscribers-mlir @llvm/pr-subscribers-mlir-openacc Author: Razvan Lupusoru (razvanlupusoru) ChangesAdd utilities in OpenACCUtilsTiling.h/.cpp to support tiling transformations on acc.loop operations:
These utilities prepare for the ACCLoopTiling pass which handles the OpenACC loop tile directive. Patch is 31.65 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/171490.diff 5 Files Affected:
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
new file mode 100644
index 0000000000000..3152526cc0582
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
@@ -0,0 +1,83 @@
+//===- OpenACCUtilsTiling.h - OpenACC Loop Tiling Utilities -----*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains utility functions for tiling OpenACC loops.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_OPENACC_OPENACCUTILSTILING_H_
+#define MLIR_DIALECT_OPENACC_OPENACCUTILSTILING_H_
+
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/PatternMatch.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace mlir {
+namespace acc {
+
+/// Uncollapse tile loops with multiple IVs and collapseCount < tileCount.
+/// This is used to prepare loops for tiling when the collapse count is less
+/// than the tile count.
+///
+/// \param origLoop The original loop operation to uncollapse.
+/// \param tileCount The number of tile dimensions.
+/// \param collapseCount The collapse count from the original loop.
+/// \param rewriter The rewriter to use for modifications.
+/// \return A vector of uncollapsed loop operations.
+llvm::SmallVector<mlir::acc::LoopOp>
+uncollapseLoops(mlir::acc::LoopOp origLoop, unsigned tileCount,
+ unsigned collapseCount, mlir::RewriterBase &rewriter);
+
+/// Tile ACC loops according to the given tile sizes.
+///
+/// Tiling a 2-level nested loop will create two 'tile' loops containing two
+/// 'element' loops. The transformation looks like:
+///
+/// Before Tiling:
+/// \code
+/// #pragma acc loop tile(tile_size1, tile_size2)
+/// for (i = lb1; i < ub1; i += step1) { // original loop
+/// for (j = lb2; j < ub2; j += step2) {
+/// a[i,j] = i + j;
+/// }
+/// }
+/// \endcode
+///
+/// After Tiling:
+/// \code
+/// for (i = lb1; i < ub1; i += (step1 * tile_size1)) { // tile loop 1
+/// for (j = lb2; j < ub2; j += (step2 * tile_size2)) { // tile loop 2
+/// for (ii = i; ii < min(ub1, (step1 * tile_size1) + i); ii += step1) {
+/// // element loop 1
+/// for (jj = j; jj < min(ub2, (step2 * tile_size2) + j); jj += step2)
+/// { // element loop 2
+/// a[ii,jj] = i + j;
+/// }
+/// }
+/// }
+/// }
+/// \endcode
+///
+/// Unknown tile sizes (represented as -1 in OpenACC for `tile(*)`) are
+/// resolved to the provided default tile size.
+///
+/// \param tileLoops The loops to tile (outermost first).
+/// \param tileSizes The tile sizes for each dimension. Values of -1 are
+/// treated as unknown and resolved to defaultTileSize.
+/// \param defaultTileSize The default tile size to use for unknown (*) tiles.
+/// \param rewriter The rewriter to use for modifications.
+/// \return The outermost loop after tiling.
+mlir::acc::LoopOp tileACCLoops(llvm::SmallVector<mlir::acc::LoopOp> &tileLoops,
+ const llvm::SmallVector<mlir::Value> &tileSizes,
+ int32_t defaultTileSize,
+ mlir::RewriterBase &rewriter);
+
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_OPENACCUTILSTILING_H_
diff --git a/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt
index 68e124625921f..c3de4f7e3e282 100644
--- a/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt
+++ b/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt
@@ -1,4 +1,5 @@
add_mlir_dialect_library(MLIROpenACCUtils
+ OpenACCUtilsTiling.cpp
OpenACCUtils.cpp
ADDITIONAL_HEADER_DIRS
@@ -14,7 +15,9 @@ add_mlir_dialect_library(MLIROpenACCUtils
MLIROpenACCTypeInterfacesIncGen
LINK_LIBS PUBLIC
+ MLIRArithDialect
MLIROpenACCDialect
MLIRIR
MLIRSupport
+ MLIRTransformUtils
)
diff --git a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
new file mode 100644
index 0000000000000..f939ec1c58cfd
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
@@ -0,0 +1,313 @@
+//===- OpenACCUtilsTiling.cpp - OpenACC Loop Tiling Utilities -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains utility functions for tiling OpenACC loops.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/OpenACCUtilsTiling.h"
+
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/Dialect/Utils/StaticValueUtils.h"
+#include "mlir/Transforms/RegionUtils.h"
+
+// Resolve unknown tile sizes (represented as -1 for tile(*)) to the default.
+static mlir::Value resolveUnknownTileSize(mlir::Value tileSize,
+ int32_t defaultTileSize,
+ mlir::RewriterBase &rewriter,
+ mlir::Location loc) {
+ auto constVal = mlir::getConstantIntValue(tileSize);
+ if (constVal && *constVal < 0) {
+ return mlir::arith::ConstantOp::create(
+ rewriter, loc, rewriter.getI32Type(),
+ rewriter.getI32IntegerAttr(defaultTileSize));
+ }
+ return tileSize;
+}
+
+// Remove vector/worker attributes from loop
+static void removeWorkerVectorFromLoop(mlir::acc::LoopOp loop) {
+ if (loop.hasVector() || loop.getVectorValue()) {
+ loop.removeVectorAttr();
+ loop.removeVectorOperandsDeviceTypeAttr();
+ } else if (loop.hasWorker() || loop.getWorkerValue()) {
+ loop.removeWorkerAttr();
+ loop.removeWorkerNumOperandsDeviceTypeAttr();
+ }
+}
+
+// Create a new ACC loop with new steps, lb, ub from original loop
+static mlir::acc::LoopOp
+createACCLoopFromOriginal(mlir::acc::LoopOp origLoop,
+ mlir::RewriterBase &rewriter, mlir::ValueRange lb,
+ mlir::ValueRange ub, mlir::ValueRange step,
+ mlir::DenseBoolArrayAttr inclusiveUBAttr,
+ mlir::acc::CombinedConstructsTypeAttr combinedAttr,
+ mlir::Location loc, bool preserveCollapse) {
+ mlir::ArrayAttr collapseAttr = mlir::ArrayAttr{};
+ mlir::ArrayAttr collapseDeviceTypeAttr = mlir::ArrayAttr{};
+ if (preserveCollapse) {
+ collapseAttr = origLoop.getCollapseAttr();
+ collapseDeviceTypeAttr = origLoop.getCollapseDeviceTypeAttr();
+ }
+ auto newLoop = mlir::acc::LoopOp::create(
+ rewriter, loc, origLoop->getResultTypes(), lb, ub, step, inclusiveUBAttr,
+ collapseAttr, collapseDeviceTypeAttr, origLoop.getGangOperands(),
+ origLoop.getGangOperandsArgTypeAttr(),
+ origLoop.getGangOperandsSegmentsAttr(),
+ origLoop.getGangOperandsDeviceTypeAttr(), origLoop.getWorkerNumOperands(),
+ origLoop.getWorkerNumOperandsDeviceTypeAttr(),
+ origLoop.getVectorOperands(), origLoop.getVectorOperandsDeviceTypeAttr(),
+ origLoop.getSeqAttr(), origLoop.getIndependentAttr(),
+ origLoop.getAuto_Attr(), origLoop.getGangAttr(), origLoop.getWorkerAttr(),
+ origLoop.getVectorAttr(), mlir::ValueRange{}, mlir::DenseI32ArrayAttr{},
+ mlir::ArrayAttr{}, origLoop.getCacheOperands(),
+ origLoop.getPrivateOperands(), origLoop.getFirstprivateOperands(),
+ origLoop.getReductionOperands(), combinedAttr);
+ return newLoop;
+}
+
+// Create inner loop inside input loop
+static mlir::acc::LoopOp
+createInnerLoop(mlir::acc::LoopOp inputLoop, mlir::RewriterBase &rewriter,
+ mlir::ValueRange lb, mlir::ValueRange ub, mlir::ValueRange step,
+ mlir::DenseBoolArrayAttr inclusiveUBAttr, mlir::Location loc) {
+ mlir::acc::LoopOp elementLoop = createACCLoopFromOriginal(
+ inputLoop, rewriter, lb, ub, step, inclusiveUBAttr,
+ mlir::acc::CombinedConstructsTypeAttr{}, loc, /*preserveCollapse*/ false);
+
+ // Remove gang/worker attributes from inner loops
+ rewriter.startOpModification(elementLoop);
+ if (inputLoop.hasGang() ||
+ inputLoop.getGangValue(mlir::acc::GangArgType::Num) ||
+ inputLoop.getGangValue(mlir::acc::GangArgType::Dim) ||
+ inputLoop.getGangValue(mlir::acc::GangArgType::Static)) {
+ elementLoop.removeGangAttr();
+ elementLoop.removeGangOperandsArgTypeAttr();
+ elementLoop.removeGangOperandsSegmentsAttr();
+ elementLoop.removeGangOperandsDeviceTypeAttr();
+ }
+ if (inputLoop.hasVector() || inputLoop.getVectorValue()) {
+ elementLoop.removeWorkerAttr();
+ elementLoop.removeWorkerNumOperandsDeviceTypeAttr();
+ }
+ rewriter.finalizeOpModification(elementLoop);
+
+ // Create empty block in elementLoop and add IV argument
+ mlir::Block *blk = rewriter.createBlock(&elementLoop.getRegion(),
+ elementLoop.getRegion().begin());
+ rewriter.setInsertionPointToEnd(blk);
+ mlir::acc::YieldOp::create(rewriter, loc);
+ elementLoop.getBody().addArgument(
+ inputLoop.getBody().getArgument(0).getType(), loc);
+
+ return elementLoop;
+}
+
+// Move ops from source to target Loop and replace uses of IVs
+static void moveOpsAndReplaceIVs(mlir::acc::LoopOp sourceLoop,
+ mlir::acc::LoopOp targetLoop,
+ llvm::ArrayRef<mlir::Value> newIVs,
+ llvm::ArrayRef<mlir::Value> origIVs,
+ size_t nOps, mlir::RewriterBase &rewriter) {
+ // Move ops from source to target loop [begin, begin + nOps - 1)
+ mlir::Block::iterator begin = sourceLoop.getBody().begin();
+ targetLoop.getBody().getOperations().splice(
+ targetLoop.getBody().getOperations().begin(),
+ sourceLoop.getBody().getOperations(), begin, std::next(begin, nOps - 1));
+
+ // Replace uses of origIV with newIV
+ for (auto [i, newIV] : llvm::enumerate(newIVs))
+ mlir::replaceAllUsesInRegionWith(origIVs[i], newIV, targetLoop.getRegion());
+}
+
+mlir::acc::LoopOp
+mlir::acc::tileACCLoops(llvm::SmallVector<mlir::acc::LoopOp> &tileLoops,
+ const llvm::SmallVector<mlir::Value> &tileSizes,
+ int32_t defaultTileSize, mlir::RewriterBase &rewriter) {
+ // Tile collapsed and/or nested loops
+ mlir::acc::LoopOp outerLoop = tileLoops[0];
+ const mlir::Location loc = outerLoop.getLoc();
+
+ // Resolve unknown tile sizes (tile(*) represented as -1)
+ llvm::SmallVector<mlir::Value> resolvedTileSizes;
+ rewriter.setInsertionPoint(outerLoop);
+ for (mlir::Value tileSize : tileSizes) {
+ resolvedTileSizes.push_back(
+ resolveUnknownTileSize(tileSize, defaultTileSize, rewriter, loc));
+ }
+
+ mlir::acc::LoopOp innerLoop = tileLoops[tileLoops.size() - 1];
+ llvm::SmallVector<mlir::Value, 3> origIVs;
+ llvm::SmallVector<mlir::Value, 3> origSteps;
+ llvm::SmallVector<mlir::Value, 3> origUBs;
+ llvm::SmallVector<mlir::Value, 3> newSteps;
+ llvm::SmallVector<mlir::Value, 3> newUBs;
+ llvm::SmallVector<mlir::Value, 3> newIVs;
+ size_t nOps = innerLoop.getBody().getOperations().size();
+
+ // Extract original inclusiveUBs
+ llvm::SmallVector<bool> inclusiveUBs;
+ for (auto tileLoop : tileLoops) {
+ for (auto [j, step] : llvm::enumerate(tileLoop.getStep())) {
+ // inclusiveUBs are present on the IR from Fortran frontend for DO loops
+ // but might not be present from other frontends (python)
+ // So check if it exists
+ if (tileLoop.getInclusiveUpperboundAttr())
+ inclusiveUBs.push_back(
+ tileLoop.getInclusiveUpperboundAttr().asArrayRef()[j]);
+ else
+ inclusiveUBs.push_back(false);
+ }
+ }
+
+ // Extract original ivs, UBs, steps, and calculate new steps
+ rewriter.setInsertionPoint(outerLoop);
+ for (auto [i, tileLoop] : llvm::enumerate(tileLoops)) {
+ for (auto arg : tileLoop.getBody().getArguments())
+ origIVs.push_back(arg);
+ for (auto ub : tileLoop.getUpperbound())
+ origUBs.push_back(ub);
+
+ llvm::SmallVector<mlir::Value, 3> currentLoopSteps;
+ for (auto [j, step] : llvm::enumerate(tileLoop.getStep())) {
+ origSteps.push_back(step);
+ if (i + j >= resolvedTileSizes.size()) {
+ currentLoopSteps.push_back(step);
+ } else {
+ mlir::Value tileSize = resolvedTileSizes[i + j];
+ auto newLoopStep =
+ mlir::arith::MulIOp::create(rewriter, loc, step, tileSize);
+ currentLoopSteps.push_back(newLoopStep);
+ newSteps.push_back(newLoopStep);
+ }
+ }
+
+ rewriter.startOpModification(tileLoop);
+ tileLoop.getStepMutable().clear();
+ tileLoop.getStepMutable().append(currentLoopSteps);
+ rewriter.finalizeOpModification(tileLoop);
+ }
+
+ // Calculate new upper bounds for element loops
+ for (size_t i = 0; i < newSteps.size(); i++) {
+ rewriter.setInsertionPoint(innerLoop.getBody().getTerminator());
+ // UpperBound: min(origUB, origIV+(originalStep*tile_size))
+ auto stepped =
+ mlir::arith::AddIOp::create(rewriter, loc, origIVs[i], newSteps[i]);
+ mlir::Value newUB = stepped;
+ if (inclusiveUBs[i]) {
+ // Handle InclusiveUB
+ // UpperBound: min(origUB, origIV+(originalStep*tile_size - 1))
+ auto c1 = mlir::arith::ConstantOp::create(
+ rewriter, loc, newSteps[i].getType(),
+ rewriter.getIntegerAttr(newSteps[i].getType(), 1));
+ newUB = mlir::arith::SubIOp::create(rewriter, loc, stepped, c1);
+ }
+ newUBs.push_back(
+ mlir::arith::MinSIOp::create(rewriter, loc, origUBs[i], newUB));
+ }
+
+ // Create and insert nested elementLoopOps before terminator of outer loopOp
+ mlir::acc::LoopOp currentLoop = innerLoop;
+ for (size_t i = 0; i < resolvedTileSizes.size(); i++) {
+ rewriter.setInsertionPoint(currentLoop.getBody().getTerminator());
+ mlir::DenseBoolArrayAttr inclusiveUBAttr = mlir::DenseBoolArrayAttr{};
+ if (inclusiveUBs[i])
+ inclusiveUBAttr = rewriter.getDenseBoolArrayAttr({true});
+
+ mlir::acc::LoopOp elementLoop =
+ createInnerLoop(innerLoop, rewriter, mlir::ValueRange{origIVs[i]},
+ mlir::ValueRange{newUBs[i]},
+ mlir::ValueRange{origSteps[i]}, inclusiveUBAttr, loc);
+
+ // Remove vector/worker attributes from inner element loops except
+ // outermost element loop
+ if (i > 0) {
+ rewriter.startOpModification(elementLoop);
+ removeWorkerVectorFromLoop(elementLoop);
+ rewriter.finalizeOpModification(elementLoop);
+ }
+ newIVs.push_back(elementLoop.getBody().getArgument(0));
+ currentLoop = elementLoop;
+ }
+
+ // Remove vector/worker attributes from outer tile loops
+ for (auto tileLoop : tileLoops) {
+ rewriter.startOpModification(tileLoop);
+ removeWorkerVectorFromLoop(tileLoop);
+ rewriter.finalizeOpModification(tileLoop);
+ }
+
+ // Move ops from inner tile loop to inner element loop and replace IV uses
+ moveOpsAndReplaceIVs(innerLoop, currentLoop, newIVs, origIVs, nOps, rewriter);
+
+ return outerLoop;
+}
+
+llvm::SmallVector<mlir::acc::LoopOp>
+mlir::acc::uncollapseLoops(mlir::acc::LoopOp origLoop, unsigned tileCount,
+ unsigned collapseCount,
+ mlir::RewriterBase &rewriter) {
+ llvm::SmallVector<mlir::acc::LoopOp, 3> newLoops;
+ llvm::SmallVector<mlir::Value, 3> newIVs;
+ mlir::Location loc = origLoop.getLoc();
+ llvm::SmallVector<bool> newInclusiveUBs;
+ llvm::SmallVector<mlir::Value, 3> lbs, ubs, steps;
+ for (unsigned i = 0; i < collapseCount; i++) {
+ // inclusiveUpperbound attribute might not be set, default to false
+ bool inclusiveUB = false;
+ if (origLoop.getInclusiveUpperboundAttr())
+ inclusiveUB = origLoop.getInclusiveUpperboundAttr().asArrayRef()[i];
+ newInclusiveUBs.push_back(inclusiveUB);
+ lbs.push_back(origLoop.getLowerbound()[i]);
+ ubs.push_back(origLoop.getUpperbound()[i]);
+ steps.push_back(origLoop.getStep()[i]);
+ }
+ mlir::acc::LoopOp outerLoop = createACCLoopFromOriginal(
+ origLoop, rewriter, lbs, ubs, steps,
+ rewriter.getDenseBoolArrayAttr(newInclusiveUBs),
+ origLoop.getCombinedAttr(), loc, /*preserveCollapse*/ true);
+ mlir::Block *blk = rewriter.createBlock(&outerLoop.getRegion(),
+ outerLoop.getRegion().begin());
+ rewriter.setInsertionPointToEnd(blk);
+ mlir::acc::YieldOp::create(rewriter, loc);
+ for (unsigned i = 0; i < collapseCount; i++) {
+ outerLoop.getBody().addArgument(origLoop.getBody().getArgument(i).getType(),
+ loc);
+ newIVs.push_back(outerLoop.getBody().getArgument(i));
+ }
+ newLoops.push_back(outerLoop);
+
+ mlir::acc::LoopOp currentLoopOp = outerLoop;
+ for (unsigned i = collapseCount; i < tileCount; i++) {
+ rewriter.setInsertionPoint(currentLoopOp.getBody().getTerminator());
+ bool inclusiveUB = false;
+ if (origLoop.getInclusiveUpperboundAttr())
+ inclusiveUB = origLoop.getInclusiveUpperboundAttr().asArrayRef()[i];
+ mlir::DenseBoolArrayAttr inclusiveUBAttr =
+ rewriter.getDenseBoolArrayAttr({inclusiveUB});
+ mlir::acc::LoopOp innerLoop = createInnerLoop(
+ origLoop, rewriter, mlir::ValueRange{origLoop.getLowerbound()[i]},
+ mlir::ValueRange{origLoop.getUpperbound()[i]},
+ mlir::ValueRange{origLoop.getStep()[i]}, inclusiveUBAttr, loc);
+ newIVs.push_back(innerLoop.getBody().getArgument(0));
+ newLoops.push_back(innerLoop);
+ currentLoopOp = innerLoop;
+ }
+ // Move ops from origLoop to innermost loop and replace uses of IVs
+ size_t nOps = origLoop.getBody().getOperations().size();
+ llvm::SmallVector<mlir::Value, 3> origIVs;
+ for (auto arg : origLoop.getBody().getArguments())
+ origIVs.push_back(arg);
+ moveOpsAndReplaceIVs(origLoop, currentLoopOp, newIVs, origIVs, nOps,
+ rewriter);
+
+ return newLoops;
+}
diff --git a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
index c8c2bb96b0539..060c8b8d2679d 100644
--- a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
+++ b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
@@ -2,6 +2,7 @@ add_mlir_unittest(MLIROpenACCTests
OpenACCOpsTest.cpp
OpenACCOpsInterfacesTest.cpp
OpenACCUtilsTest.cpp
+ OpenACCUtilsTilingTest.cpp
)
mlir_target_link_libraries(MLIROpenACCTests
PRIVATE
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
new file mode 100644
index 0000000000000..287af9fafd5b7
--- /dev/null
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
@@ -0,0 +1,349 @@
+//===- OpenACCUtilsTilingTest.cpp - Unit tests for loop tiling utilities --===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/OpenACCUtilsTiling.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/OwningOpRef.h"
+#include "gtest/gtest.h"
+
+using namespace mlir;
+using namespace mlir::acc;
+
+//===----------------------------------------------------------------------===//
+// Test Fixture
+//===----------------------------------------------------------------------===//
+
+class OpenACCUtilsTilingTest : public ::testing::Test {
+protected:
+ OpenACCUtilsTilingTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ context.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
+ memref::MemRefDialect, func::FuncDialect>(...
[truncated]
|
Member
|
@llvm/pr-subscribers-openacc Author: Razvan Lupusoru (razvanlupusoru) ChangesAdd utilities in OpenACCUtilsTiling.h/.cpp to support tiling transformations on acc.loop operations:
These utilities prepare for the ACCLoopTiling pass which handles the OpenACC loop tile directive. Patch is 31.65 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/171490.diff 5 Files Affected:
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
new file mode 100644
index 0000000000000..3152526cc0582
--- /dev/null
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
@@ -0,0 +1,83 @@
+//===- OpenACCUtilsTiling.h - OpenACC Loop Tiling Utilities -----*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains utility functions for tiling OpenACC loops.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef MLIR_DIALECT_OPENACC_OPENACCUTILSTILING_H_
+#define MLIR_DIALECT_OPENACC_OPENACCUTILSTILING_H_
+
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/PatternMatch.h"
+#include "llvm/ADT/SmallVector.h"
+
+namespace mlir {
+namespace acc {
+
+/// Uncollapse tile loops with multiple IVs and collapseCount < tileCount.
+/// This is used to prepare loops for tiling when the collapse count is less
+/// than the tile count.
+///
+/// \param origLoop The original loop operation to uncollapse.
+/// \param tileCount The number of tile dimensions.
+/// \param collapseCount The collapse count from the original loop.
+/// \param rewriter The rewriter to use for modifications.
+/// \return A vector of uncollapsed loop operations.
+llvm::SmallVector<mlir::acc::LoopOp>
+uncollapseLoops(mlir::acc::LoopOp origLoop, unsigned tileCount,
+ unsigned collapseCount, mlir::RewriterBase &rewriter);
+
+/// Tile ACC loops according to the given tile sizes.
+///
+/// Tiling a 2-level nested loop will create two 'tile' loops containing two
+/// 'element' loops. The transformation looks like:
+///
+/// Before Tiling:
+/// \code
+/// #pragma acc loop tile(tile_size1, tile_size2)
+/// for (i = lb1; i < ub1; i += step1) { // original loop
+/// for (j = lb2; j < ub2; j += step2) {
+/// a[i,j] = i + j;
+/// }
+/// }
+/// \endcode
+///
+/// After Tiling:
+/// \code
+/// for (i = lb1; i < ub1; i += (step1 * tile_size1)) { // tile loop 1
+/// for (j = lb2; j < ub2; j += (step2 * tile_size2)) { // tile loop 2
+/// for (ii = i; ii < min(ub1, (step1 * tile_size1) + i); ii += step1) {
+/// // element loop 1
+/// for (jj = j; jj < min(ub2, (step2 * tile_size2) + j); jj += step2)
+/// { // element loop 2
+/// a[ii,jj] = i + j;
+/// }
+/// }
+/// }
+/// }
+/// \endcode
+///
+/// Unknown tile sizes (represented as -1 in OpenACC for `tile(*)`) are
+/// resolved to the provided default tile size.
+///
+/// \param tileLoops The loops to tile (outermost first).
+/// \param tileSizes The tile sizes for each dimension. Values of -1 are
+/// treated as unknown and resolved to defaultTileSize.
+/// \param defaultTileSize The default tile size to use for unknown (*) tiles.
+/// \param rewriter The rewriter to use for modifications.
+/// \return The outermost loop after tiling.
+mlir::acc::LoopOp tileACCLoops(llvm::SmallVector<mlir::acc::LoopOp> &tileLoops,
+ const llvm::SmallVector<mlir::Value> &tileSizes,
+ int32_t defaultTileSize,
+ mlir::RewriterBase &rewriter);
+
+} // namespace acc
+} // namespace mlir
+
+#endif // MLIR_DIALECT_OPENACC_OPENACCUTILSTILING_H_
diff --git a/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt b/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt
index 68e124625921f..c3de4f7e3e282 100644
--- a/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt
+++ b/mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt
@@ -1,4 +1,5 @@
add_mlir_dialect_library(MLIROpenACCUtils
+ OpenACCUtilsTiling.cpp
OpenACCUtils.cpp
ADDITIONAL_HEADER_DIRS
@@ -14,7 +15,9 @@ add_mlir_dialect_library(MLIROpenACCUtils
MLIROpenACCTypeInterfacesIncGen
LINK_LIBS PUBLIC
+ MLIRArithDialect
MLIROpenACCDialect
MLIRIR
MLIRSupport
+ MLIRTransformUtils
)
diff --git a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
new file mode 100644
index 0000000000000..f939ec1c58cfd
--- /dev/null
+++ b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
@@ -0,0 +1,313 @@
+//===- OpenACCUtilsTiling.cpp - OpenACC Loop Tiling Utilities -------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file contains utility functions for tiling OpenACC loops.
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/OpenACCUtilsTiling.h"
+
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/Dialect/Utils/StaticValueUtils.h"
+#include "mlir/Transforms/RegionUtils.h"
+
+// Resolve unknown tile sizes (represented as -1 for tile(*)) to the default.
+static mlir::Value resolveUnknownTileSize(mlir::Value tileSize,
+ int32_t defaultTileSize,
+ mlir::RewriterBase &rewriter,
+ mlir::Location loc) {
+ auto constVal = mlir::getConstantIntValue(tileSize);
+ if (constVal && *constVal < 0) {
+ return mlir::arith::ConstantOp::create(
+ rewriter, loc, rewriter.getI32Type(),
+ rewriter.getI32IntegerAttr(defaultTileSize));
+ }
+ return tileSize;
+}
+
+// Remove vector/worker attributes from loop
+static void removeWorkerVectorFromLoop(mlir::acc::LoopOp loop) {
+ if (loop.hasVector() || loop.getVectorValue()) {
+ loop.removeVectorAttr();
+ loop.removeVectorOperandsDeviceTypeAttr();
+ } else if (loop.hasWorker() || loop.getWorkerValue()) {
+ loop.removeWorkerAttr();
+ loop.removeWorkerNumOperandsDeviceTypeAttr();
+ }
+}
+
+// Create a new ACC loop with new steps, lb, ub from original loop
+static mlir::acc::LoopOp
+createACCLoopFromOriginal(mlir::acc::LoopOp origLoop,
+ mlir::RewriterBase &rewriter, mlir::ValueRange lb,
+ mlir::ValueRange ub, mlir::ValueRange step,
+ mlir::DenseBoolArrayAttr inclusiveUBAttr,
+ mlir::acc::CombinedConstructsTypeAttr combinedAttr,
+ mlir::Location loc, bool preserveCollapse) {
+ mlir::ArrayAttr collapseAttr = mlir::ArrayAttr{};
+ mlir::ArrayAttr collapseDeviceTypeAttr = mlir::ArrayAttr{};
+ if (preserveCollapse) {
+ collapseAttr = origLoop.getCollapseAttr();
+ collapseDeviceTypeAttr = origLoop.getCollapseDeviceTypeAttr();
+ }
+ auto newLoop = mlir::acc::LoopOp::create(
+ rewriter, loc, origLoop->getResultTypes(), lb, ub, step, inclusiveUBAttr,
+ collapseAttr, collapseDeviceTypeAttr, origLoop.getGangOperands(),
+ origLoop.getGangOperandsArgTypeAttr(),
+ origLoop.getGangOperandsSegmentsAttr(),
+ origLoop.getGangOperandsDeviceTypeAttr(), origLoop.getWorkerNumOperands(),
+ origLoop.getWorkerNumOperandsDeviceTypeAttr(),
+ origLoop.getVectorOperands(), origLoop.getVectorOperandsDeviceTypeAttr(),
+ origLoop.getSeqAttr(), origLoop.getIndependentAttr(),
+ origLoop.getAuto_Attr(), origLoop.getGangAttr(), origLoop.getWorkerAttr(),
+ origLoop.getVectorAttr(), mlir::ValueRange{}, mlir::DenseI32ArrayAttr{},
+ mlir::ArrayAttr{}, origLoop.getCacheOperands(),
+ origLoop.getPrivateOperands(), origLoop.getFirstprivateOperands(),
+ origLoop.getReductionOperands(), combinedAttr);
+ return newLoop;
+}
+
+// Create inner loop inside input loop
+static mlir::acc::LoopOp
+createInnerLoop(mlir::acc::LoopOp inputLoop, mlir::RewriterBase &rewriter,
+ mlir::ValueRange lb, mlir::ValueRange ub, mlir::ValueRange step,
+ mlir::DenseBoolArrayAttr inclusiveUBAttr, mlir::Location loc) {
+ mlir::acc::LoopOp elementLoop = createACCLoopFromOriginal(
+ inputLoop, rewriter, lb, ub, step, inclusiveUBAttr,
+ mlir::acc::CombinedConstructsTypeAttr{}, loc, /*preserveCollapse*/ false);
+
+ // Remove gang/worker attributes from inner loops
+ rewriter.startOpModification(elementLoop);
+ if (inputLoop.hasGang() ||
+ inputLoop.getGangValue(mlir::acc::GangArgType::Num) ||
+ inputLoop.getGangValue(mlir::acc::GangArgType::Dim) ||
+ inputLoop.getGangValue(mlir::acc::GangArgType::Static)) {
+ elementLoop.removeGangAttr();
+ elementLoop.removeGangOperandsArgTypeAttr();
+ elementLoop.removeGangOperandsSegmentsAttr();
+ elementLoop.removeGangOperandsDeviceTypeAttr();
+ }
+ if (inputLoop.hasVector() || inputLoop.getVectorValue()) {
+ elementLoop.removeWorkerAttr();
+ elementLoop.removeWorkerNumOperandsDeviceTypeAttr();
+ }
+ rewriter.finalizeOpModification(elementLoop);
+
+ // Create empty block in elementLoop and add IV argument
+ mlir::Block *blk = rewriter.createBlock(&elementLoop.getRegion(),
+ elementLoop.getRegion().begin());
+ rewriter.setInsertionPointToEnd(blk);
+ mlir::acc::YieldOp::create(rewriter, loc);
+ elementLoop.getBody().addArgument(
+ inputLoop.getBody().getArgument(0).getType(), loc);
+
+ return elementLoop;
+}
+
+// Move ops from source to target Loop and replace uses of IVs
+static void moveOpsAndReplaceIVs(mlir::acc::LoopOp sourceLoop,
+ mlir::acc::LoopOp targetLoop,
+ llvm::ArrayRef<mlir::Value> newIVs,
+ llvm::ArrayRef<mlir::Value> origIVs,
+ size_t nOps, mlir::RewriterBase &rewriter) {
+ // Move ops from source to target loop [begin, begin + nOps - 1)
+ mlir::Block::iterator begin = sourceLoop.getBody().begin();
+ targetLoop.getBody().getOperations().splice(
+ targetLoop.getBody().getOperations().begin(),
+ sourceLoop.getBody().getOperations(), begin, std::next(begin, nOps - 1));
+
+ // Replace uses of origIV with newIV
+ for (auto [i, newIV] : llvm::enumerate(newIVs))
+ mlir::replaceAllUsesInRegionWith(origIVs[i], newIV, targetLoop.getRegion());
+}
+
+mlir::acc::LoopOp
+mlir::acc::tileACCLoops(llvm::SmallVector<mlir::acc::LoopOp> &tileLoops,
+ const llvm::SmallVector<mlir::Value> &tileSizes,
+ int32_t defaultTileSize, mlir::RewriterBase &rewriter) {
+ // Tile collapsed and/or nested loops
+ mlir::acc::LoopOp outerLoop = tileLoops[0];
+ const mlir::Location loc = outerLoop.getLoc();
+
+ // Resolve unknown tile sizes (tile(*) represented as -1)
+ llvm::SmallVector<mlir::Value> resolvedTileSizes;
+ rewriter.setInsertionPoint(outerLoop);
+ for (mlir::Value tileSize : tileSizes) {
+ resolvedTileSizes.push_back(
+ resolveUnknownTileSize(tileSize, defaultTileSize, rewriter, loc));
+ }
+
+ mlir::acc::LoopOp innerLoop = tileLoops[tileLoops.size() - 1];
+ llvm::SmallVector<mlir::Value, 3> origIVs;
+ llvm::SmallVector<mlir::Value, 3> origSteps;
+ llvm::SmallVector<mlir::Value, 3> origUBs;
+ llvm::SmallVector<mlir::Value, 3> newSteps;
+ llvm::SmallVector<mlir::Value, 3> newUBs;
+ llvm::SmallVector<mlir::Value, 3> newIVs;
+ size_t nOps = innerLoop.getBody().getOperations().size();
+
+ // Extract original inclusiveUBs
+ llvm::SmallVector<bool> inclusiveUBs;
+ for (auto tileLoop : tileLoops) {
+ for (auto [j, step] : llvm::enumerate(tileLoop.getStep())) {
+ // inclusiveUBs are present on the IR from Fortran frontend for DO loops
+ // but might not be present from other frontends (python)
+ // So check if it exists
+ if (tileLoop.getInclusiveUpperboundAttr())
+ inclusiveUBs.push_back(
+ tileLoop.getInclusiveUpperboundAttr().asArrayRef()[j]);
+ else
+ inclusiveUBs.push_back(false);
+ }
+ }
+
+ // Extract original ivs, UBs, steps, and calculate new steps
+ rewriter.setInsertionPoint(outerLoop);
+ for (auto [i, tileLoop] : llvm::enumerate(tileLoops)) {
+ for (auto arg : tileLoop.getBody().getArguments())
+ origIVs.push_back(arg);
+ for (auto ub : tileLoop.getUpperbound())
+ origUBs.push_back(ub);
+
+ llvm::SmallVector<mlir::Value, 3> currentLoopSteps;
+ for (auto [j, step] : llvm::enumerate(tileLoop.getStep())) {
+ origSteps.push_back(step);
+ if (i + j >= resolvedTileSizes.size()) {
+ currentLoopSteps.push_back(step);
+ } else {
+ mlir::Value tileSize = resolvedTileSizes[i + j];
+ auto newLoopStep =
+ mlir::arith::MulIOp::create(rewriter, loc, step, tileSize);
+ currentLoopSteps.push_back(newLoopStep);
+ newSteps.push_back(newLoopStep);
+ }
+ }
+
+ rewriter.startOpModification(tileLoop);
+ tileLoop.getStepMutable().clear();
+ tileLoop.getStepMutable().append(currentLoopSteps);
+ rewriter.finalizeOpModification(tileLoop);
+ }
+
+ // Calculate new upper bounds for element loops
+ for (size_t i = 0; i < newSteps.size(); i++) {
+ rewriter.setInsertionPoint(innerLoop.getBody().getTerminator());
+ // UpperBound: min(origUB, origIV+(originalStep*tile_size))
+ auto stepped =
+ mlir::arith::AddIOp::create(rewriter, loc, origIVs[i], newSteps[i]);
+ mlir::Value newUB = stepped;
+ if (inclusiveUBs[i]) {
+ // Handle InclusiveUB
+ // UpperBound: min(origUB, origIV+(originalStep*tile_size - 1))
+ auto c1 = mlir::arith::ConstantOp::create(
+ rewriter, loc, newSteps[i].getType(),
+ rewriter.getIntegerAttr(newSteps[i].getType(), 1));
+ newUB = mlir::arith::SubIOp::create(rewriter, loc, stepped, c1);
+ }
+ newUBs.push_back(
+ mlir::arith::MinSIOp::create(rewriter, loc, origUBs[i], newUB));
+ }
+
+ // Create and insert nested elementLoopOps before terminator of outer loopOp
+ mlir::acc::LoopOp currentLoop = innerLoop;
+ for (size_t i = 0; i < resolvedTileSizes.size(); i++) {
+ rewriter.setInsertionPoint(currentLoop.getBody().getTerminator());
+ mlir::DenseBoolArrayAttr inclusiveUBAttr = mlir::DenseBoolArrayAttr{};
+ if (inclusiveUBs[i])
+ inclusiveUBAttr = rewriter.getDenseBoolArrayAttr({true});
+
+ mlir::acc::LoopOp elementLoop =
+ createInnerLoop(innerLoop, rewriter, mlir::ValueRange{origIVs[i]},
+ mlir::ValueRange{newUBs[i]},
+ mlir::ValueRange{origSteps[i]}, inclusiveUBAttr, loc);
+
+ // Remove vector/worker attributes from inner element loops except
+ // outermost element loop
+ if (i > 0) {
+ rewriter.startOpModification(elementLoop);
+ removeWorkerVectorFromLoop(elementLoop);
+ rewriter.finalizeOpModification(elementLoop);
+ }
+ newIVs.push_back(elementLoop.getBody().getArgument(0));
+ currentLoop = elementLoop;
+ }
+
+ // Remove vector/worker attributes from outer tile loops
+ for (auto tileLoop : tileLoops) {
+ rewriter.startOpModification(tileLoop);
+ removeWorkerVectorFromLoop(tileLoop);
+ rewriter.finalizeOpModification(tileLoop);
+ }
+
+ // Move ops from inner tile loop to inner element loop and replace IV uses
+ moveOpsAndReplaceIVs(innerLoop, currentLoop, newIVs, origIVs, nOps, rewriter);
+
+ return outerLoop;
+}
+
+llvm::SmallVector<mlir::acc::LoopOp>
+mlir::acc::uncollapseLoops(mlir::acc::LoopOp origLoop, unsigned tileCount,
+ unsigned collapseCount,
+ mlir::RewriterBase &rewriter) {
+ llvm::SmallVector<mlir::acc::LoopOp, 3> newLoops;
+ llvm::SmallVector<mlir::Value, 3> newIVs;
+ mlir::Location loc = origLoop.getLoc();
+ llvm::SmallVector<bool> newInclusiveUBs;
+ llvm::SmallVector<mlir::Value, 3> lbs, ubs, steps;
+ for (unsigned i = 0; i < collapseCount; i++) {
+ // inclusiveUpperbound attribute might not be set, default to false
+ bool inclusiveUB = false;
+ if (origLoop.getInclusiveUpperboundAttr())
+ inclusiveUB = origLoop.getInclusiveUpperboundAttr().asArrayRef()[i];
+ newInclusiveUBs.push_back(inclusiveUB);
+ lbs.push_back(origLoop.getLowerbound()[i]);
+ ubs.push_back(origLoop.getUpperbound()[i]);
+ steps.push_back(origLoop.getStep()[i]);
+ }
+ mlir::acc::LoopOp outerLoop = createACCLoopFromOriginal(
+ origLoop, rewriter, lbs, ubs, steps,
+ rewriter.getDenseBoolArrayAttr(newInclusiveUBs),
+ origLoop.getCombinedAttr(), loc, /*preserveCollapse*/ true);
+ mlir::Block *blk = rewriter.createBlock(&outerLoop.getRegion(),
+ outerLoop.getRegion().begin());
+ rewriter.setInsertionPointToEnd(blk);
+ mlir::acc::YieldOp::create(rewriter, loc);
+ for (unsigned i = 0; i < collapseCount; i++) {
+ outerLoop.getBody().addArgument(origLoop.getBody().getArgument(i).getType(),
+ loc);
+ newIVs.push_back(outerLoop.getBody().getArgument(i));
+ }
+ newLoops.push_back(outerLoop);
+
+ mlir::acc::LoopOp currentLoopOp = outerLoop;
+ for (unsigned i = collapseCount; i < tileCount; i++) {
+ rewriter.setInsertionPoint(currentLoopOp.getBody().getTerminator());
+ bool inclusiveUB = false;
+ if (origLoop.getInclusiveUpperboundAttr())
+ inclusiveUB = origLoop.getInclusiveUpperboundAttr().asArrayRef()[i];
+ mlir::DenseBoolArrayAttr inclusiveUBAttr =
+ rewriter.getDenseBoolArrayAttr({inclusiveUB});
+ mlir::acc::LoopOp innerLoop = createInnerLoop(
+ origLoop, rewriter, mlir::ValueRange{origLoop.getLowerbound()[i]},
+ mlir::ValueRange{origLoop.getUpperbound()[i]},
+ mlir::ValueRange{origLoop.getStep()[i]}, inclusiveUBAttr, loc);
+ newIVs.push_back(innerLoop.getBody().getArgument(0));
+ newLoops.push_back(innerLoop);
+ currentLoopOp = innerLoop;
+ }
+ // Move ops from origLoop to innermost loop and replace uses of IVs
+ size_t nOps = origLoop.getBody().getOperations().size();
+ llvm::SmallVector<mlir::Value, 3> origIVs;
+ for (auto arg : origLoop.getBody().getArguments())
+ origIVs.push_back(arg);
+ moveOpsAndReplaceIVs(origLoop, currentLoopOp, newIVs, origIVs, nOps,
+ rewriter);
+
+ return newLoops;
+}
diff --git a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
index c8c2bb96b0539..060c8b8d2679d 100644
--- a/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
+++ b/mlir/unittests/Dialect/OpenACC/CMakeLists.txt
@@ -2,6 +2,7 @@ add_mlir_unittest(MLIROpenACCTests
OpenACCOpsTest.cpp
OpenACCOpsInterfacesTest.cpp
OpenACCUtilsTest.cpp
+ OpenACCUtilsTilingTest.cpp
)
mlir_target_link_libraries(MLIROpenACCTests
PRIVATE
diff --git a/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
new file mode 100644
index 0000000000000..287af9fafd5b7
--- /dev/null
+++ b/mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp
@@ -0,0 +1,349 @@
+//===- OpenACCUtilsTilingTest.cpp - Unit tests for loop tiling utilities --===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "mlir/Dialect/OpenACC/OpenACCUtilsTiling.h"
+#include "mlir/Dialect/Arith/IR/Arith.h"
+#include "mlir/Dialect/Func/IR/FuncOps.h"
+#include "mlir/Dialect/MemRef/IR/MemRef.h"
+#include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/BuiltinOps.h"
+#include "mlir/IR/BuiltinTypes.h"
+#include "mlir/IR/MLIRContext.h"
+#include "mlir/IR/OwningOpRef.h"
+#include "gtest/gtest.h"
+
+using namespace mlir;
+using namespace mlir::acc;
+
+//===----------------------------------------------------------------------===//
+// Test Fixture
+//===----------------------------------------------------------------------===//
+
+class OpenACCUtilsTilingTest : public ::testing::Test {
+protected:
+ OpenACCUtilsTilingTest() : b(&context), loc(UnknownLoc::get(&context)) {
+ context.loadDialect<acc::OpenACCDialect, arith::ArithDialect,
+ memref::MemRefDialect, func::FuncDialect>(...
[truncated]
|
clementval
reviewed
Dec 9, 2025
|
✅ With the latest revision this PR passed the C/C++ code formatter. |
clementval
reviewed
Dec 9, 2025
added 2 commits
December 9, 2025 11:01
clementval
reviewed
Dec 9, 2025
clementval
reviewed
Dec 9, 2025
clementval
approved these changes
Dec 9, 2025
Collaborator
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/203/builds/31383 Here is the relevant piece of the build log for the reference |
Collaborator
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/205/builds/30173 Here is the relevant piece of the build log for the reference |
Contributor
Author
|
The buildbot failures should get resolved with: #171546 |
Collaborator
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/117/builds/15627 Here is the relevant piece of the build log for the reference |
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
5 participants
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Add utilities in OpenACCUtilsTiling.h/.cpp to support tiling transformations on acc.loop operations:
These utilities prepare for the ACCLoopTiling pass which handles the OpenACC loop tile directive.