[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:

I'd slightly prefer keeping this as a single `st.async` op. From the codegen side, splitting doesn't buy us anything, the generator now has to decide which op to emit based on the address space, which just pushes the complexity from the verifier into every producer of the op. 

A slightly larger verifier is a non-issue to me; that's exactly where this kind of constraint checking belongs.

Plus it's confusing to see multiple OPs while PTX has only one. It's hard to find. 

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