[mlir][acc] Introduce acc loop tiling pass (#171692)

This pass implements the OpenACC loop tiling transformation for acc.loop
operations that have the tile clause (OpenACC 3.4 spec, section 2.9.8).

The tile clause specifies that the iterations of the associated loops
should be divided into tiles (rectangular blocks). The 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).

For example, tiling a 2-level nested loop with tile(T1, T2):
```
  // Before tiling:
  acc.loop tile(T1, T2) control(%i, %j) = ...

  // After tiling:
  acc.loop control(%i) step (s1*T1) {        // tile loop 1
    acc.loop control(%j) step (s2*T2) {      // tile loop 2
      acc.loop control(%ii) = (%i) to (min(ub1, %i+s1*T1)) {
        acc.loop control(%jj) = (%j) to (min(ub2, %j+s2*T2)) {
          // loop body using %ii, %jj
        }
      }
    }
  }
```

Key features:
- Handles constant tile sizes and wildcard tile sizes ('*') which use a
configurable default tile size
- Properly handles collapsed loops with tile counts exceeding collapse
count by uncollapsing loops before tiling
- Distributes gang/worker/vector attributes appropriately: gang -> tile
loops, vector -> element loops
- Validates that tile size types are not wider than loop IV types
- Emits optimization remarks for tiling decisions

Three test files are added:
- acc-loop-tiling.mlir: Tests single and nested loop tiling with
constant tile sizes, unknown tile sizes (*), and loops with collapse
attributes
- acc-loop-tiling-invalid.mlir: Tests error diagnostic when tile size
type is wider than the loop IV type
- acc-loop-tiling-remarks.mlir: Tests optimization remarks emitted for
tiling decisions including default tile size selection

Co-authored-by: Vijay Kandiah <[email protected]>

Author: razvanlupusoru

mlir/include/mlir/Dialect/OpenACC/Analysis/OpenACCSupport.h

@@ -50,8 +50,10 @@
 #ifndef MLIR_DIALECT_OPENACC_ANALYSIS_OPENACCSUPPORT_H
 #define MLIR_DIALECT_OPENACC_ANALYSIS_OPENACCSUPPORT_H
 
+#include "mlir/IR/Remarks.h"
 #include "mlir/IR/Value.h"
 #include "mlir/Pass/AnalysisManager.h"
+#include "llvm/ADT/StringRef.h"
 #include <memory>
 #include <string>
 
@@ -62,6 +64,8 @@ namespace acc {
 enum class RecipeKind : uint32_t;
 bool isValidSymbolUse(Operation *user, SymbolRefAttr symbol,
                       Operation **definingOpPtr);
+remark::detail::InFlightRemark emitRemark(Operation *op, const Twine &message,
+                                          llvm::StringRef category);
 
 namespace detail {
 /// This class contains internal trait classes used by OpenACCSupport.
@@ -82,6 +86,14 @@ struct OpenACCSupportTraits {
     // Used to report a case that is not supported by the implementation.
     virtual InFlightDiagnostic emitNYI(Location loc, const Twine &message) = 0;
 
+    // Used to emit an OpenACC remark. The category is optional and is used to
+    // either capture the pass name or pipeline phase when the remark is
+    // emitted. When not provided, in the default implementation, the category
+    // is "openacc".
+    virtual remark::detail::InFlightRemark
+    emitRemark(Operation *op, const Twine &message,
+               llvm::StringRef category) = 0;
+
     /// Check if a symbol use is valid for use in an OpenACC region.
     virtual bool isValidSymbolUse(Operation *user, SymbolRefAttr symbol,
                                   Operation **definingOpPtr) = 0;
@@ -101,13 +113,22 @@ struct OpenACCSupportTraits {
                         Operation **>;
 
   template <typename ImplT, typename... Args>
+
   using isValidValueUse_t =
       decltype(std::declval<ImplT>().isValidValueUse(std::declval<Args>()...));
 
   template <typename ImplT>
   using has_isValidValueUse =
       llvm::is_detected<isValidValueUse_t, ImplT, Value, Region &>;
 
+  template <typename ImplT, typename... Args>
+  using emitRemark_t =
+      decltype(std::declval<ImplT>().emitRemark(std::declval<Args>()...));
+
+  template <typename ImplT>
+  using has_emitRemark = llvm::is_detected<emitRemark_t, ImplT, Operation *,
+                                           const Twine &, llvm::StringRef>;
+
   /// This class wraps a concrete OpenACCSupport implementation and forwards
   /// interface calls to it. This provides type erasure, allowing different
   /// implementation types to be used interchangeably without inheritance.
@@ -131,6 +152,15 @@ struct OpenACCSupportTraits {
       return impl.emitNYI(loc, message);
     }
 
+    remark::detail::InFlightRemark emitRemark(Operation *op,
+                                              const Twine &message,
+                                              llvm::StringRef category) final {
+      if constexpr (has_emitRemark<ImplT>::value)
+        return impl.emitRemark(op, message, category);
+      else
+        return acc::emitRemark(op, message, category);
+    }
+
     bool isValidSymbolUse(Operation *user, SymbolRefAttr symbol,
                           Operation **definingOpPtr) final {
       if constexpr (has_isValidSymbolUse<ImplT>::value)
@@ -198,6 +228,17 @@ class OpenACCSupport {
   ///         unsupported case.
   InFlightDiagnostic emitNYI(Location loc, const Twine &message);
 
+  /// Emit an OpenACC remark.
+  ///
+  /// \param op The operation to emit the remark for.
+  /// \param message The remark message.
+  /// \param category Optional category for the remark. Defaults to "openacc".
+  /// \return An in-flight remark object that can be used to append
+  ///         additional information to the remark.
+  remark::detail::InFlightRemark
+  emitRemark(Operation *op, const Twine &message,
+             llvm::StringRef category = "openacc");
+
   /// Check if a symbol use is valid for use in an OpenACC region.
   ///
   /// \param user The operation using the symbol.

mlir/include/mlir/Dialect/OpenACC/OpenACCUtils.h

@@ -10,7 +10,9 @@
 #define MLIR_DIALECT_OPENACC_OPENACCUTILS_H_
 
 #include "mlir/Dialect/OpenACC/OpenACC.h"
+#include "mlir/IR/Remarks.h"
 #include "llvm/ADT/SmallVector.h"
+#include "llvm/ADT/StringRef.h"
 
 namespace mlir {
 class DominanceInfo;
@@ -81,6 +83,17 @@ getDominatingDataClauses(mlir::Operation *computeConstructOp,
                          mlir::DominanceInfo &domInfo,
                          mlir::PostDominanceInfo &postDomInfo);
 
+/// Emit an OpenACC remark for the given operation with the given message.
+///
+/// \param op The operation to emit the remark for.
+/// \param message The remark message.
+/// \param category Optional category for the remark. Defaults to "openacc".
+/// \return An in-flight remark object that can be used to append
+///         additional information to the remark.
+remark::detail::InFlightRemark emitRemark(mlir::Operation *op,
+                                          const llvm::Twine &message,
+                                          llvm::StringRef category = "openacc");
+
 } // namespace acc
 } // namespace mlir
 

mlir/include/mlir/Dialect/OpenACC/Transforms/Passes.td

@@ -152,4 +152,46 @@ def ACCLegalizeSerial : Pass<"acc-legalize-serial", "mlir::func::FuncOp"> {
       "mlir::arith::ArithDialect"];
 }
 
+
+def ACCLoopTiling : Pass<"acc-loop-tiling", "mlir::func::FuncOp"> {
+  let summary = "Tile OpenACC loops with tile clauses";
+  let description = [{
+    This pass implements loop tiling transformations for OpenACC loops that
+    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
+
+    Example transformation:
+    ```
+    // Before:
+    #pragma acc loop tile(32, 32)
+    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] = ...
+    ```
+
+    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::acc::OpenACCDialect",
+      "mlir::arith::ArithDialect"];
+  let options = [
+    Option<"defaultTileSize", "default-tile-size", "int32_t", "32",
+           "Default tile size to use for wildcard ('*') tile sizes">
+  ];
+}
+
 #endif // MLIR_DIALECT_OPENACC_TRANSFORMS_PASSES

mlir/lib/Dialect/OpenACC/Analysis/OpenACCSupport.cpp

@@ -41,6 +41,14 @@ InFlightDiagnostic OpenACCSupport::emitNYI(Location loc, const Twine &message) {
   return mlir::emitError(loc, "not yet implemented: " + message);
 }
 
+remark::detail::InFlightRemark
+OpenACCSupport::emitRemark(Operation *op, const Twine &message,
+                           llvm::StringRef category) {
+  if (impl)
+    return impl->emitRemark(op, message, category);
+  return acc::emitRemark(op, message, category);
+}
+
 bool OpenACCSupport::isValidSymbolUse(Operation *user, SymbolRefAttr symbol,
                                       Operation **definingOpPtr) {
   if (impl)