[Mlir-commits] [mlir] [MLIR][NVVM] Add support for mapa MLIR Ops (PR #124514)

[Guray Ozen]

================
@@ -2512,6 +2512,43 @@ def NVVM_WgmmaMmaAsyncOp : NVVM_Op<"wgmma.mma_async",
   }];
 }
 
+//===----------------------------------------------------------------------===//
+// NVVM Mapa Ops
+//===----------------------------------------------------------------------===//
+
+def NVVM_MapaOp : NVVM_IntrOp<"mapa", [], 1> {
+  let results = (outs LLVM_PointerGeneric:$res);
+  let arguments = (ins LLVM_PointerGeneric:$a, I32:$b);
+
+  let description = [{
+    Maps the generic address pointing to a shared memory variable in the 
+    target CTA. Source `a` and `res` are registers containing generic 
+    addresses pointing to shared memory. 
+    `b` is a 32-bit integer operand representing the rank of the target CTA. 
+    [For more information, see PTX ISA]
+    (https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-mapa)
+  }];
+    
+  let assemblyFormat ="$a`,` $b attr-dict";
+}
+
+def NVVM_MapaSharedClusterOp : NVVM_IntrOp<"mapa.shared.cluster", [], 1> {
----------------
grypp wrote:

PR looks great, thanks for implementing.
I am wondering if we should combine the OPs or not. They look quite similar. 

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