[Mlir-commits] [flang] [mlir] [mlir][acc] Apply tile sizes to the innermost loop first (PR #237355)
Slava Zakharin
llvmlistbot at llvm.org
Mon Nov 2 11:04:43 PST 2026
https://github.com/vzakhari updated https://github.com/llvm/llvm-project/pull/237355
>From af1e118052693e0479f495c7533adb96eda3dd72 Mon Sep 17 00:00:00 2001
From: Slava Zakharin <szakharin at nvidia.com>
Date: Fri, 30 Oct 2026 11:20:49 -0700
Subject: [PATCH 1/2] [mlir][acc] Apply tile sizes to the innermost loop first
The OpenACC tile clause lists its sizes innermost loop first (OpenACC 3.4,
section 2.9.8). In
!$acc loop tile(128, 1)
do k = 1, nk
do i = 1, ni
the i loop is tiled by 128 and the k loop by 1. Frontends store the sizes
on acc.loop in clause order and list the induction variables outermost
first, but tileACCLoops paired the size at position k with induction
variable k, so the nest above was tiled by 128 along k and by 1 along i.
The results stay correct, but the tile is transposed: for an array indexed
(i, k), the vector lanes of a tile access memory with a stride instead of
contiguously.
Pair the size at position k with induction variable n - 1 - k in
tileACCLoops and in the pass's tile size type check, and document the
operand order on acc.loop. Since the type check now compares each size with
the induction variable it tiles, a nest with an integer(8) outer and a
default integer inner DO variable is now accepted with tile(8, 4_8) and
reported as not yet implemented with tile(4_8, 8), the reverse of before.
Also drop the claim in the pass description that collapsed loops are
handled: since #211651, a tile clause combined with a collapse clause on
the same loop is reported as not yet implemented.
Assisted-by: Cursor
---
flang/test/Lower/OpenACC/acc-loop.f90 | 6 +-
.../mlir/Dialect/OpenACC/OpenACCOps.td | 15 +++-
.../mlir/Dialect/OpenACC/OpenACCUtilsTiling.h | 19 +++--
.../mlir/Dialect/OpenACC/Transforms/Passes.td | 23 +++---
.../OpenACC/Transforms/ACCLoopTiling.cpp | 40 +++++-----
.../OpenACC/Utils/OpenACCUtilsTiling.cpp | 8 +-
.../OpenACC/acc-loop-tiling-invalid.mlir | 21 ++++++
.../test/Dialect/OpenACC/acc-loop-tiling.mlir | 75 ++++++++++++++++++-
8 files changed, 159 insertions(+), 48 deletions(-)
diff --git a/flang/test/Lower/OpenACC/acc-loop.f90 b/flang/test/Lower/OpenACC/acc-loop.f90
index 8354c663eb788f..de0cd1a888d220 100644
--- a/flang/test/Lower/OpenACC/acc-loop.f90
+++ b/flang/test/Lower/OpenACC/acc-loop.f90
@@ -234,15 +234,17 @@ program acc_loop
! CHECK: acc.yield
! CHECK-NEXT: } inclusiveUpperbound(array<i1: true>) independent
- !$acc loop tile(2, 2)
+ !$acc loop tile(2, 4)
DO i = 1, n
DO j = 1, n
c(i, j) = d(i, j)
END DO
END DO
+! The tile sizes keep their clause order; the induction variables are listed
+! outermost first.
! CHECK: [[TILESIZE1:%.*]] = arith.constant 2 : i32
-! CHECK: [[TILESIZE2:%.*]] = arith.constant 2 : i32
+! CHECK: [[TILESIZE2:%.*]] = arith.constant 4 : i32
! CHECK: acc.loop {{.*}} tile({[[TILESIZE1]] : i32, [[TILESIZE2]] : i32}) control(%arg0 : i32, %arg1 : i32) = (%{{.*}} : i32, i32) to (%{{.*}} : i32, i32) step (%{{.*}} : i32, i32) {
! CHECK: acc.yield
! CHECK-NEXT: } inclusiveUpperbound(array<i1: true, true>) independent
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
index 47457abd250656..6f71ddffcd7c37 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCOps.td
@@ -2687,6 +2687,13 @@ def OpenACC_LoopOp
as Loop Invariant Code Motion breaks the property requested by the
clause on the loop constructs.
+ The `tile` operands hold the sizes of the `tile` clause in source order.
+ As defined by the OpenACC specification (version 3.4, section 2.9.8), the
+ first size applies to the innermost associated loop and the last size to
+ the outermost one: with `n` sizes, the size at position `k` applies to
+ induction variable `n - 1 - k`. A size of `-1` represents `*`, an
+ implementation-defined tile size.
+
Each `acc.loop` holds private and reduction operands which are the
ssa values from the corresponding `acc.private` or `acc.reduction`
operations. Additionally, firstprivate operands are supported to
@@ -2826,10 +2833,11 @@ def OpenACC_LoopOp
/// device_type.
bool hasWorker(mlir::acc::DeviceType deviceType);
- /// Return the values of the tile clause if present.
+ /// Return the values of the tile clause if present, in clause order: the
+ /// first value is the tile size of the innermost associated loop.
mlir::Operation::operand_range getTileValues();
/// Return the values of the tile clause for the given device_type if
- /// present.
+ /// present, in clause order.
mlir::Operation::operand_range
getTileValues(mlir::acc::DeviceType deviceType);
@@ -2866,7 +2874,8 @@ def OpenACC_LoopOp
void setCollapseForDeviceTypes(MLIRContext *, llvm::ArrayRef<DeviceType>,
llvm::APInt);
// Sets the tile values for this 'loop' for a set of DeviceTypes. All of the
- // values should be integral constants, with the '*' represented as a '-1'.
+ // values should be integral constants in clause order, with the '*'
+ // represented as a '-1'.
void setTileForDeviceTypes(MLIRContext *, llvm::ArrayRef<DeviceType>,
mlir::ValueRange);
diff --git a/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
index b24ab07c37c5cf..eaff9b2cb4708d 100644
--- a/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
+++ b/mlir/include/mlir/Dialect/OpenACC/OpenACCUtilsTiling.h
@@ -19,7 +19,6 @@
namespace mlir {
namespace acc {
-
/// Tile a single fused acc.loop that carries all associated induction
/// variables (one IV per tile dimension).
///
@@ -45,14 +44,16 @@ namespace acc {
/// }
/// \endcode
///
-/// After Tiling (each group is one multi-IV loop over all tiled IVs):
+/// After Tiling (each group is one multi-IV loop over all tiled IVs). As the
+/// OpenACC specification defines, the first tile size applies to the innermost
+/// loop:
/// \code
/// // tile group
-/// for (i = lb1; i < ub1; i += (step1 * tile_size1),
-/// j = lb2; j < ub2; j += (step2 * tile_size2)) {
+/// for (i = lb1; i < ub1; i += (step1 * tile_size2),
+/// j = lb2; j < ub2; j += (step2 * tile_size1)) {
/// // element group
-/// for (ii = i; ii < min(ub1, (step1 * tile_size1) + i); ii += step1,
-/// jj = j; jj < min(ub2, (step2 * tile_size2) + j); jj += step2) {
+/// for (ii = i; ii < min(ub1, (step1 * tile_size2) + i); ii += step1,
+/// jj = j; jj < min(ub2, (step2 * tile_size1) + j); jj += step2) {
/// a[ii,jj] = i + j;
/// }
/// }
@@ -66,8 +67,10 @@ namespace acc {
/// resolved to the provided default tile size.
///
/// \param tileLoop The fused loop to tile.
-/// \param tileSizes The tile sizes for each tiled dimension. Values of -1 are
-/// treated as unknown and resolved to defaultTileSize.
+/// \param tileSizes The tile sizes in the order of the tile clause: with `n`
+/// sizes, the first `n` induction variables are tiled and the size at
+/// position `k` applies to induction variable `n - 1 - k`. 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 tile group loop that is modified in place.
diff --git a/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td b/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
index 202028994003dd..3cbb6f702b588a 100644
--- a/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
+++ b/mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td
@@ -268,30 +268,33 @@ def ACCLoopTiling : Pass<"acc-loop-tiling", "mlir::func::FuncOp"> {
have tile clauses. The pass transforms loops with `tile(size1, size2, ...)`
clauses into tiled loop nests.
- For a 2-level nested loop with tile(T1, T2), the transformation produces:
- - Outer tile loops that iterate over tiles
- - Inner element loops that iterate within each tile
+ The transformation produces two loops, each carrying all the induction
+ variables of the original loop:
+ - A tile group loop that iterates over tiles
+ - An element group loop that iterates within each tile
+
+ The tile clause lists its sizes innermost loop first (OpenACC 3.4,
+ section 2.9.8). In the example below, 32 tiles the `j` loop and 4 tiles
+ the `i` loop.
Example transformation:
```
// Before:
- #pragma acc loop tile(32, 32)
+ #pragma acc loop tile(32, 4)
for (i = 0; i < N; i++)
for (j = 0; j < M; j++)
A[i][j] = ...
// After:
- for (i = 0; i < N; i += 32) // tile loop 1
- for (j = 0; j < M; j += 32) // tile loop 2
- for (ii = i; ii < min(N, i+32); ii++) // element loop 1
- for (jj = j; jj < min(M, j+32); jj++) // element loop 2
- A[ii][jj] = ...
+ for (i = 0; i < N; i += 4, j = 0; j < M; j += 32) // tile group
+ for (ii = i; ii < min(N, i+4); ii++,
+ jj = j; jj < min(M, j+32); jj++) // element group
+ A[ii][jj] = ...
```
The pass handles:
- Constant tile sizes
- Wildcard tile sizes ('*') which use a default tile size
- - Collapsed loops with tile counts exceeding collapse count
- Proper handling of loop attributes (gang, worker, vector)
}];
let dependentDialects = ["mlir::arith::ArithDialect"];
diff --git a/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp b/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp
index 9f71b5d77cd558..dd4e0d0218d6b4 100644
--- a/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp
+++ b/mlir/lib/Dialect/OpenACC/Transforms/ACCLoopTiling.cpp
@@ -12,26 +12,26 @@
// Overview:
// ---------
// The tile clause specifies that the iterations of the associated loops should
-// be divided into tiles (rectangular blocks). This pass transforms a single
-// or nested acc.loop with tile clauses into a structure of "tile loops"
-// (iterating over tiles) containing "element loops" (iterating within tiles).
+// be divided into tiles (rectangular blocks). This pass transforms an acc.loop
+// with a tile clause into a "tile group" loop (iterating over tiles) containing
+// an "element group" loop (iterating within a tile). Each of them carries all
+// the induction variables of the original loop.
//
-// For example, tiling a 2-level nested loop with tile(T1, T2) produces:
+// The tile clause lists its sizes innermost loop first, so the first size
+// applies to the last induction variable. For example, tiling a 2-level loop
+// nest with tile(T1, T2) produces:
//
// // Before tiling:
-// acc.loop tile(T1, T2) control(%i, %j) = (lb1, lb2) to (ub1, ub2) step (s1,
-// s2)
+// acc.loop tile(T1, T2) control(%i, %j) = (lb1, lb2) to (ub1, ub2)
+// step (s1, s2)
//
// // After tiling:
-// acc.loop control(%i) = (lb1) to (ub1) step (s1*T1) { // tile loop 1
-// acc.loop control(%j) = (lb2) to (ub2) step (s2*T2) { // tile loop 2
-// acc.loop control(%ii) = (%i) to (min(ub1, %i+s1*T1)) step (s1) { //
-// element 1
-// acc.loop control(%jj) = (%j) to (min(ub2, %j+s2*T2)) step (s2) { //
-// element 2
-// // loop body using %ii, %jj
-// }
-// }
+// acc.loop control(%i, %j) = (lb1, lb2) to (ub1, ub2)
+// step (s1*T2, s2*T1) { // tile group
+// acc.loop control(%ii, %jj) = (%i, %j)
+// to (min(ub1, %i+s1*T2), min(ub2, %j+s2*T1))
+// step (s1, s2) { // element group
+// // loop body using %ii, %jj
// }
// }
//
@@ -92,15 +92,17 @@ struct ACCLoopTilingImpl : public OpRewritePattern<acc::LoopOp> {
// Check that tile size types are not narrower than IV types.
// We only check when both types are IntegerType. For IndexType, the width
// is target-dependent and the casting utility will handle it correctly.
+ // The tile sizes are in clause order, so the last one belongs to the
+ // outermost IV.
LogicalResult checkTileSizeTypes(acc::LoopOp loop,
ArrayRef<Value> tileSizes) const {
auto ivTypes = loop.getBody().getArgumentTypes();
- for (size_t i = 0; i < tileSizes.size() && i < ivTypes.size(); ++i) {
- Type tileType = tileSizes[i].getType();
- Type ivType = ivTypes[i];
+ for (auto [tileSize, ivType] :
+ llvm::zip(llvm::reverse(tileSizes), ivTypes)) {
+ Type tileType = tileSize.getType();
// Skip unknown tile sizes (will be created with correct type)
- auto constVal = getConstantIntValue(tileSizes[i]);
+ auto constVal = getConstantIntValue(tileSize);
if (constVal && *constVal < 0)
continue;
diff --git a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
index 24242f86dd4f28..3026831e804444 100644
--- a/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
+++ b/mlir/lib/Dialect/OpenACC/Utils/OpenACCUtilsTiling.cpp
@@ -179,14 +179,16 @@ mlir::acc::tileACCLoops(mlir::acc::LoopOp tileLoop,
}
// Scale each tiled dimension's step by its tile size to form the tile group
- // loop steps.
+ // loop steps. The tile sizes are in clause order, so the first one belongs
+ // to the innermost tiled IV.
rewriter.setInsertionPoint(tileLoop);
llvm::SmallVector<mlir::Value, 3> scaledSteps;
llvm::SmallVector<mlir::Value, 3> tileLoopSteps;
for (unsigned i = 0; i < numIVs; ++i) {
if (i < tileCount) {
- mlir::Value tileSize = resolveAndCastTileSize(
- tileSizes[i], defaultTileSize, origSteps[i].getType(), rewriter, loc);
+ mlir::Value tileSize =
+ resolveAndCastTileSize(tileSizes[tileCount - 1 - i], defaultTileSize,
+ origSteps[i].getType(), rewriter, loc);
mlir::Value scaled =
mlir::arith::MulIOp::create(rewriter, loc, origSteps[i], tileSize);
scaledSteps.push_back(scaled);
diff --git a/mlir/test/Dialect/OpenACC/acc-loop-tiling-invalid.mlir b/mlir/test/Dialect/OpenACC/acc-loop-tiling-invalid.mlir
index 4214b367c42cc8..2e90bb0bdfe8ba 100644
--- a/mlir/test/Dialect/OpenACC/acc-loop-tiling-invalid.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-loop-tiling-invalid.mlir
@@ -16,6 +16,27 @@ func.func @tile_wider_than_iv(%arg0: memref<100xf32>) {
// -----
+// The first tile size applies to the innermost loop, so the i64 size below
+// tiles the i32 IV %j.
+
+func.func @tile_wider_than_inner_iv(%arg0: memref<100x50xf32>) {
+ %c0_i64 = arith.constant 0 : i64
+ %c100_i64 = arith.constant 100 : i64
+ %c1_i64 = arith.constant 1 : i64
+ %c4_i64 = arith.constant 4 : i64
+ %c0_i32 = arith.constant 0 : i32
+ %c50_i32 = arith.constant 50 : i32
+ %c1_i32 = arith.constant 1 : i32
+ %c8_i32 = arith.constant 8 : i32
+ // expected-error @+1 {{not yet implemented: tile size type (i64) is wider than loop IV type (i32)}}
+ acc.loop tile({%c4_i64 : i64, %c8_i32 : i32}) control(%i : i64, %j : i32) = (%c0_i64, %c0_i32 : i64, i32) to (%c100_i64, %c50_i32 : i64, i32) step (%c1_i64, %c1_i32 : i64, i32) {
+ acc.yield
+ } independent
+ return
+}
+
+// -----
+
// A tile clause combined with a collapse clause on the same loop is not
// supported: the two clauses describe conflicting loop-association counts. The
// pass must diagnose it rather than silently drop the collapse clause.
diff --git a/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir b/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir
index 6db2060124e9a9..0e50784ec3e33e 100644
--- a/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir
@@ -34,6 +34,8 @@ func.func @single_loop_tile(%arg0: memref<10xf32>) {
// Test 2-level nested loop tiling with tile(4, 8)
// Produces one multi-IV tile-group loop (gang) wrapping one multi-IV
// element-group loop (vector), each carrying both tiled induction variables.
+// The first tile size applies to the innermost loop: %j is tiled by 4 and %i
+// by 8.
// CHECK-LABEL: func.func @nested_loop_tile
// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
@@ -43,9 +45,11 @@ func.func @single_loop_tile(%arg0: memref<10xf32>) {
// CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index
// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : index
// Tile-group loop with gang; steps scaled by the tile sizes.
-// CHECK: acc.loop gang control(%[[I:.*]] : index, %[[J:.*]] : index) = (%[[C0]], %[[C0]] : index, index) to (%[[C100]], %[[C50]] : index, index) step (%[[C4]], %[[C8]] : index, index) {
-// CHECK: %[[MUB0:.*]] = arith.minsi
-// CHECK: %[[MUB1:.*]] = arith.minsi
+// CHECK: acc.loop gang control(%[[I:.*]] : index, %[[J:.*]] : index) = (%[[C0]], %[[C0]] : index, index) to (%[[C100]], %[[C50]] : index, index) step (%[[C8]], %[[C4]] : index, index) {
+// CHECK: %[[E0:.*]] = arith.addi %[[I]], %[[C8]] : index
+// CHECK: %[[MUB0:.*]] = arith.minsi %[[E0]], %[[C100]] : index
+// CHECK: %[[E1:.*]] = arith.addi %[[J]], %[[C4]] : index
+// CHECK: %[[MUB1:.*]] = arith.minsi %[[E1]], %[[C50]] : index
// Element-group loop with vector.
// CHECK: acc.loop vector control(%{{.*}} : index, %{{.*}} : index) = (%[[I]], %[[J]] : index, index) to (%[[MUB0]], %[[MUB1]] : index, index) step (%[[C1]], %[[C1]] : index, index) {
// CHECK: acc.yield
@@ -69,6 +73,71 @@ func.func @nested_loop_tile(%arg0: memref<100x50xf32>) {
return
}
+// Test 3-level loop tiling with tile(2, 4, 8): %k is tiled by 2, %j by 4 and
+// %i by 8.
+
+// CHECK-LABEL: func.func @three_dim_tile
+// CHECK-DAG: %[[C0:.*]] = arith.constant 0 : index
+// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK-DAG: %[[C2:.*]] = arith.constant 2 : index
+// CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index
+// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : index
+// CHECK-DAG: %[[C10:.*]] = arith.constant 10 : index
+// CHECK-DAG: %[[C20:.*]] = arith.constant 20 : index
+// CHECK-DAG: %[[C30:.*]] = arith.constant 30 : index
+// CHECK: acc.loop gang control(%[[I:.*]] : index, %[[J:.*]] : index, %[[K:.*]] : index) = (%[[C0]], %[[C0]], %[[C0]] : index, index, index) to (%[[C10]], %[[C20]], %[[C30]] : index, index, index) step (%[[C8]], %[[C4]], %[[C2]] : index, index, index) {
+// CHECK: %[[E0:.*]] = arith.addi %[[I]], %[[C8]] : index
+// CHECK: %[[MUB0:.*]] = arith.minsi %[[E0]], %[[C10]] : index
+// CHECK: %[[E1:.*]] = arith.addi %[[J]], %[[C4]] : index
+// CHECK: %[[MUB1:.*]] = arith.minsi %[[E1]], %[[C20]] : index
+// CHECK: %[[E2:.*]] = arith.addi %[[K]], %[[C2]] : index
+// CHECK: %[[MUB2:.*]] = arith.minsi %[[E2]], %[[C30]] : index
+// CHECK: acc.loop vector control(%{{.*}} : index, %{{.*}} : index, %{{.*}} : index) = (%[[I]], %[[J]], %[[K]] : index, index, index) to (%[[MUB0]], %[[MUB1]], %[[MUB2]] : index, index, index) step (%[[C1]], %[[C1]], %[[C1]] : index, index, index) {
+func.func @three_dim_tile(%arg0: memref<10x20x30xf32>) {
+ %c0 = arith.constant 0 : index
+ %c10 = arith.constant 10 : index
+ %c20 = arith.constant 20 : index
+ %c30 = arith.constant 30 : index
+ %c1 = arith.constant 1 : index
+ %c2 = arith.constant 2 : index
+ %c4 = arith.constant 4 : index
+ %c8 = arith.constant 8 : index
+ %cst = arith.constant 1.0 : f32
+ acc.loop gang vector tile({%c2 : index, %c4 : index, %c8 : index}) control(%i : index, %j : index, %k : index) = (%c0, %c0, %c0 : index, index, index) to (%c10, %c20, %c30 : index, index, index) step (%c1, %c1, %c1 : index, index, index) {
+ memref.store %cst, %arg0[%i, %j, %k] : memref<10x20x30xf32>
+ acc.yield
+ } independent
+ return
+}
+
+// Test tile sizes of different widths: the i32 size tiles the innermost i32
+// IV and the i64 size tiles the outermost i64 IV, so no size is wider than
+// the IV it tiles.
+
+// CHECK-LABEL: func.func @mixed_width_tile
+// CHECK-DAG: %[[C4:.*]] = arith.constant 4 : i64
+// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : i32
+// CHECK: acc.loop gang control(%{{.*}} : i64, %{{.*}} : i32) = {{.*}} step (%[[C4]], %[[C8]] : i64, i32) {
+// CHECK: acc.loop vector control(%{{.*}} : i64, %{{.*}} : i32)
+func.func @mixed_width_tile(%arg0: memref<100x50xf32>) {
+ %c0_i64 = arith.constant 0 : i64
+ %c100_i64 = arith.constant 100 : i64
+ %c1_i64 = arith.constant 1 : i64
+ %c4_i64 = arith.constant 4 : i64
+ %c0_i32 = arith.constant 0 : i32
+ %c50_i32 = arith.constant 50 : i32
+ %c1_i32 = arith.constant 1 : i32
+ %c8_i32 = arith.constant 8 : i32
+ %cst = arith.constant 1.0 : f32
+ acc.loop gang vector tile({%c8_i32 : i32, %c4_i64 : i64}) control(%i : i64, %j : i32) = (%c0_i64, %c0_i32 : i64, i32) to (%c100_i64, %c50_i32 : i64, i32) step (%c1_i64, %c1_i32 : i64, i32) {
+ %ii = arith.index_cast %i : i64 to index
+ %jj = arith.index_cast %j : i32 to index
+ memref.store %cst, %arg0[%ii, %jj] : memref<100x50xf32>
+ acc.yield
+ } independent
+ return
+}
+
// Regression test: a loop with GANG(STATIC: N) combined with a multi-dim
// TILE clause used to crash with an assertion failure inside
// LoopOp::getGangValue() when a leftover gang operand was copied onto a
>From ac72f9372a7f0fe988489283a2cafc3d16017202 Mon Sep 17 00:00:00 2001
From: Slava Zakharin <szakharin at nvidia.com>
Date: Mon, 2 Nov 2026 10:42:01 -0800
Subject: [PATCH 2/2] [mlir][acc][NFC] Test tiling a loop nest with inclusive
upper bounds
Flang lowers every DO loop nest to an acc.loop with inclusive upper bounds,
but the ACCLoopTiling tests only tiled single loops with inclusive bounds and
never checked the inclusiveUpperbound attribute of the generated loops. Add a
2-level case that checks the element bound of each induction variable and
that both the tile group and the element group keep the inclusive bounds.
Assisted-by: Cursor
---
.../test/Dialect/OpenACC/acc-loop-tiling.mlir | 32 +++++++++++++++++++
1 file changed, 32 insertions(+)
diff --git a/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir b/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir
index 0e50784ec3e33e..a06692433d9dfd 100644
--- a/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir
+++ b/mlir/test/Dialect/OpenACC/acc-loop-tiling.mlir
@@ -351,6 +351,38 @@ func.func @ascending_inclusive_tile(%arg0: memref<128xf32>, %lb: index, %ub: ind
return
}
+// Test 2-level loop tiling with inclusive upper bounds, the form of a Fortran
+// DO loop nest: each element bound subtracts one from its tile's edge, and
+// both the tile group and the element group keep the inclusive upper bounds.
+
+// CHECK-LABEL: func.func @nested_inclusive_tile
+// CHECK-DAG: %[[C1:.*]] = arith.constant 1 : index
+// CHECK-DAG: %[[C4:.*]] = arith.constant 4 : index
+// CHECK-DAG: %[[C8:.*]] = arith.constant 8 : index
+// CHECK: acc.loop gang control(%[[I:.*]] : index, %[[J:.*]] : index) = (%[[C1]], %[[C1]] : index, index) to (%[[UB0:.*]], %[[UB1:.*]] : index, index) step (%[[C8]], %[[C4]] : index, index) {
+// CHECK: %[[STEPPED0:.*]] = arith.addi %[[I]], %[[C8]] : index
+// CHECK: %[[EDGE0:.*]] = arith.subi %[[STEPPED0]], %{{.*}} : index
+// CHECK: %[[MUB0:.*]] = arith.minsi %[[UB0]], %[[EDGE0]] : index
+// CHECK: %[[STEPPED1:.*]] = arith.addi %[[J]], %[[C4]] : index
+// CHECK: %[[EDGE1:.*]] = arith.subi %[[STEPPED1]], %{{.*}} : index
+// CHECK: %[[MUB1:.*]] = arith.minsi %[[UB1]], %[[EDGE1]] : index
+// CHECK: acc.loop vector control(%{{.*}} : index, %{{.*}} : index) = (%[[I]], %[[J]] : index, index) to (%[[MUB0]], %[[MUB1]] : index, index) step (%[[C1]], %[[C1]] : index, index) {
+// CHECK: acc.yield
+// CHECK: } inclusiveUpperbound(array<i1: true, true>) independent
+// CHECK: acc.yield
+// CHECK: } inclusiveUpperbound(array<i1: true, true>) independent
+func.func @nested_inclusive_tile(%arg0: memref<?x?xf32>, %ub0: index, %ub1: index) {
+ %c1 = arith.constant 1 : index
+ %c4 = arith.constant 4 : index
+ %c8 = arith.constant 8 : index
+ %cst = arith.constant 1.0 : f32
+ acc.loop gang vector tile({%c4 : index, %c8 : index}) control(%i : index, %j : index) = (%c1, %c1 : index, index) to (%ub0, %ub1 : index, index) step (%c1, %c1 : index, index) {
+ memref.store %cst, %arg0[%i, %j] : memref<?x?xf32>
+ acc.yield
+ } inclusiveUpperbound(array<i1: true, true>) independent
+ return
+}
+
// Test a step whose sign is only known at run time: both clamps are computed
// and selected on the sign.
More information about the Mlir-commits
mailing list