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

[Guray Ozen]

================
@@ -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");
----------------
grypp wrote:

We also hearing complaints about discoverability. For example, PTX has a single `st.async` instruction, but NVVM has multipe OPs with different names. It makes hard to find the op if you don't know nvvm.




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