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

[Srinivasa Ravi]

================
@@ -3662,6 +3679,40 @@ DivFOp::getIntrinsicIDAndArgs(Operation &op, LLVM::ModuleTranslation &mt,
           {mt.lookupValue(thisOp.getLhs()), mt.lookupValue(thisOp.getRhs())}};
 }
 
+mlir::NVVM::IDArgPair AsyncStoreGlobalOp::getIntrinsicIDAndArgs(
+    Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+  using IDArgPair = mlir::NVVM::IDArgPair;
+  auto thisOp = cast<NVVM::AsyncStoreGlobalOp>(op);
+  mlir::NVVM::MemScopeKind scope = thisOp.getScope();
+  bool isMmio = thisOp.getMmio();
+
+  llvm::Value *addr = mt.lookupValue(thisOp.getAddr());
+  llvm::Value *value = mt.lookupValue(thisOp.getValue());
+  llvm::Value *isMultimem = builder.getInt1(thisOp.getMultimem());
+
+  if (scope == MemScopeKind::SYS) {
+    return isMmio ? IDArgPair(llvm::Intrinsic::nvvm_st_async_mmio_sys,
+                              {addr, value})
+                  : IDArgPair(llvm::Intrinsic::nvvm_st_async_sys,
+                              {addr, value, isMultimem});
+  } else if (scope == MemScopeKind::GPU) {
+    return NVVM::IDArgPair(llvm::Intrinsic::nvvm_st_async_gpu,
+                           {addr, value, isMultimem});
+  }
+  llvm_unreachable("unsupported async store scope for global address space");
+}
+
+mlir::NVVM::IDArgPair AsyncStoreSharedOp::getIntrinsicIDAndArgs(
+    Operation &op, LLVM::ModuleTranslation &mt, llvm::IRBuilderBase &builder) {
+  auto thisOp = cast<NVVM::AsyncStoreSharedOp>(op);
+
+  llvm::Value *addr = mt.lookupValue(thisOp.getAddr());
+  llvm::Value *value = mt.lookupValue(thisOp.getValue());
+  llvm::Value *mbarrier = mt.lookupValue(thisOp.getMbarrier());
+
+  return {llvm::Intrinsic::nvvm_st_async, {addr, value, mbarrier}};
----------------
Wolfram70 wrote:

Moved to using `LLVM_IntrOp` for this since it is a direct Op-intrinsic mapping in the latest revision, thanks!

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