[Mlir-commits] [mlir] [MLIR][NVVM] Update TMA tensor prefetch Op (PR #153464)

[Guray Ozen]

================
@@ -2253,6 +2253,56 @@ def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 // NVVM TMA Ops
 //===----------------------------------------------------------------------===//
 
+// List of modes supported for TMA Load and Prefetch Ops
+def TMALoadModeTile   : I32EnumAttrCase<"TILE", 0, "tile">;
+def TMALoadModeIm2Col : I32EnumAttrCase<"IM2COL", 1, "im2col">;
+def TMALoadModeIm2ColW : I32EnumAttrCase<"IM2COL_W", 2, "im2col_w">;
+def TMALoadModeIm2ColW128 : I32EnumAttrCase<"IM2COL_W_128", 3, "im2col_w_128">;
+def TMALoadModeTileGather4 : I32EnumAttrCase<"TILE_GATHER4", 4, "tile_gather4">;
----------------
grypp wrote:

So we are implementing the load mode:

 ```
.load_mode = { .tile, .tile::gather4, .im2col, .im2col::w, .im2col::w::128 }
 ```

 `tile::` and `im2col::` are implemented as main modes, while `gather4` and `w` are their sub-details.

 Is it possible to implement this in the NVVM dialect?


https://github.com/llvm/llvm-project/pull/153464