[Mlir-commits] [mlir] [mlir][NVVM] Add support for few more fence Ops (PR #170251)

[Pradeep Kumar]

================
@@ -411,6 +411,39 @@ getTcgen05StIntrinsicID(mlir::NVVM::Tcgen05LdStShape shape, uint32_t num) {
   llvm_unreachable("unhandled tcgen05.st lowering");
 }
 
+static llvm::Intrinsic::ID
+getFenceProxyID(NVVM::ProxyKind kind, std::optional<NVVM::SharedSpace> space) {
+  switch (kind) {
+  case NVVM::ProxyKind::alias:
+    return llvm::Intrinsic::nvvm_fence_proxy_alias;
+  case NVVM::ProxyKind::async:
+    return llvm::Intrinsic::nvvm_fence_proxy_async;
+  case NVVM::ProxyKind::async_global:
+    return llvm::Intrinsic::nvvm_fence_proxy_async_global;
+  case NVVM::ProxyKind::async_shared:
+    if (*space == NVVM::SharedSpace::shared_cta)
+      return llvm::Intrinsic::nvvm_fence_proxy_async_shared_cta;
+    else
+      return llvm::Intrinsic::nvvm_fence_proxy_async_shared_cluster;
+  default:
+    llvm_unreachable("unsupported proxy kind");
+  }
+}
+
+static llvm::Intrinsic::ID
+getFenceProxySyncRestrictID(NVVM::MemOrderKind order) {
+  switch (order) {
+  case NVVM::MemOrderKind::ACQUIRE:
+    return llvm::Intrinsic::
+        nvvm_fence_proxy_async_generic_acquire_sync_restrict_space_cluster_scope_cluster;
+  case NVVM::MemOrderKind::RELEASE:
+    return llvm::Intrinsic::
+        nvvm_fence_proxy_async_generic_release_sync_restrict_space_cta_scope_cluster;
+  default:
+    llvm_unreachable("unsupported memory semantics");
----------------
schwarzschild-radius wrote:

Thanks, Updated code to the above pattern

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