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

[Durgadoss R]

================
@@ -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
----------------
durga4github wrote:

else not required

+
Can a ternary simplify things here?

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