[Mlir-commits] [mlir] [mlir][nvvm] Introduce `elect.sync` Op (PR #68323)

[Guray Ozen]

================
@@ -574,6 +574,26 @@ def NVVM_SyncWarpOp :
 }
 
 
+def NVVM_ElectSyncOp : NVVM_Op<"elect.sync", 
+                  [DeclareOpInterfaceMethods<BasicPtxBuilderOpInterface>]>
+{  
+  let results = (outs I1:$pred);
+  let assemblyFormat = "attr-dict `->` type(results)";  
+  let extraClassDefinition = [{        
+    std::string $cppClass::getPtx() { 
+      return std::string(
+        "{                                  \n"
+        ".reg .u32 rx;                      \n"
+        ".reg .pred px;                     \n"
+        " mov.u32 %0, 0;                    \n"
+        "    elect.sync rx | px, 0xFFFFFFFF;\n"
----------------
grypp wrote:

>So it would be similar to using LLVM intrinsics that may or may not be available for a particular target.

Exactly, it is similar to using LLVM intrinsic. NVVM was designed to map 1to1 to LLVM intrinsic. 

FWIW, generating PTX like I do here is relatively new in MLIR. We noticed that LLVM doesn't support all of the PTX (also predicates). Therefore, I've recently implemented the `BasicPtxBuilderOpInterface` that generates `inline assembly` with the given PTX. This is what I use here. 

> Speaking of intrinsics, it may make sense to make elect.sync an LLVM intrinsic. We'll eventually need it there in any case.

Yes it does makes sense. We would like to generate LLVM intrinsic whenever it is possible. 

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