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

[Artem Belevich]

================
@@ -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"
----------------
Artem-B wrote:

Thank you for the explanation. So it would be similar to using LLVM intrinsics that may or may not be available for a particular target.

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

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