[Mlir-commits] [mlir] [mlir][acc] Add loop tiling utilities for OpenACC (PR #171490)
llvmlistbot at llvm.org
llvmlistbot at llvm.org
Tue Dec 9 10:59:10 PST 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-openacc
Author: Razvan Lupusoru (razvanlupusoru)
<details>
<summary>Changes</summary>
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.
---
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:
- (added) mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h (+83)
- (modified) mlir/lib/Dialect/OpenACC/Utils/CMakeLists.txt (+3)
- (added) mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp (+313)
- (modified) mlir/unittests/Dialect/OpenACC/CMakeLists.txt (+1)
- (added) mlir/unittests/Dialect/OpenACC/OpenACCUtilsTilingTest.cpp (+349)
``````````diff
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]
``````````
</details>
https://github.com/llvm/llvm-project/pull/171490
More information about the Mlir-commits
mailing list