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

[Guray Ozen]

================
@@ -3662,6 +3706,42 @@ DivFOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
           {mt.lookupValue(thisOp.getLhs()), mt.lookupValue(thisOp.getRhs())}};
 }
 
+mlir::NVVM::IDArgPair
+AsyncStoreOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
+                                    llvm::IRBuilderBase &builder) {
+  auto thisOp = cast<NVVM::AsyncStoreOp>(op);
+  NVVM::NVVMMemorySpace addrSpace = static_cast<NVVM::NVVMMemorySpace>(
+      llvm::cast<LLVM::LLVMPointerType>(thisOp.getAddr().getType())
+          .getAddressSpace());
+  mlir::NVVM::AsyncStoreScope scope = thisOp.getScope();
+
+  llvm::Value *addr = mt.lookupValue(thisOp.getAddr());
+  llvm::Value *value = mt.lookupValue(thisOp.getValue());
+  llvm::Value *mbarrier =
+      thisOp.getMbarrier() ? mt.lookupValue(thisOp.getMbarrier()) : nullptr;
+  llvm::Value *isMultimem =
+      thisOp.getIsMultimem() ? builder.getInt1(true) : builder.getInt1(false);
+
+  switch (addrSpace) {
+  case NVVMMemorySpace::Global:
+    if (scope == AsyncStoreScope::SYS) {
+      if (thisOp.getIsMmio())
+        return {llvm::Intrinsic::nvvm_st_async_mmio_sys, {addr, value}};
+      else
+        return {llvm::Intrinsic::nvvm_st_async_sys, {addr, value, isMultimem}};
+    } else if (scope == AsyncStoreScope::GPU) {
+      return {llvm::Intrinsic::nvvm_st_async_gpu, {addr, value, isMultimem}};
+    }
+    llvm_unreachable("unsupported async store scope for global address space");
+  case NVVMMemorySpace::SharedCluster:
+    return {llvm::Intrinsic::nvvm_st_async, {addr, value, mbarrier}};
+  default:
+    llvm_unreachable("unsupported address space");
+  }
+
+  return {llvm::Intrinsic::not_intrinsic, {}};
----------------
grypp wrote:

use `llvm_unreachables`

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