[Mlir-commits] [mlir] [MLIR][NVVM] Add asynchronous store Op (PR #210931)

[Durgadoss R]

================
@@ -668,6 +668,50 @@ LogicalResult BulkStoreOp::verify() {
   return success();
 }
 
+LogicalResult AsyncStoreOp::verify() {
+  unsigned addrSpace =
+      llvm::cast<LLVM::LLVMPointerType>(getAddr().getType()).getAddressSpace();
+  NVVM::AsyncStoreScope scope = getScope();
+  mlir::Type valueType = getValue().getType();
+  bool isMmio = getIsMmio();
+  bool isMultimem = getIsMultimem();
+
+  if (addrSpace == NVVMMemorySpace::Global) {
+    if (getMbarrier())
+      return emitOpError("mbarrier is not supported for global address space");
+
+    if (valueType.isInteger(128))
+      return emitOpError("only 8, 16, 32, and 64 bit values are supported for "
+                         "global address space");
+
+    if (scope == AsyncStoreScope::NONE)
+      return emitOpError(
+          "scope must be set for async store to global address space");
+
+    if (isMmio && scope != AsyncStoreScope::SYS)
+      return emitOpError("mmio is only supported for SYS scope");
+
+    if (isMmio && isMultimem)
+      return emitOpError("multimem is not supported for mmio");
+  }
+
+  if (addrSpace == NVVMMemorySpace::SharedCluster) {
+    if (valueType.isInteger(8) || valueType.isInteger(16))
+      return emitOpError("only 32, 64, and 128 bit values are supported for "
+                         "shared cluster address space");
+
+    if (isMultimem || isMmio)
+      return emitOpError("multimem and mmio are not supported for shared "
+                         "cluster address space");
----------------
durga4github wrote:

I would use the TMA Ops (where we have separate Ops for TMA-load and TMA-store) that perform async-data-copy as an analogy for this Op and not the `prefetch` family.

It seems we cannot implement a combined Op without updating the existing' MemAttr` or adding a new one. This seems unnecessary since we can solve it by cleanly splitting the Op.


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