[mlir][nvvm] Fix mov.u32 to mov.pred (#70027)

This PR fixes the incorrect `mov` instruction in PTX. We actually move a
predicate here, not u32, so the correct instruction should be
`mov.pred`.

Author: grypp

mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td

@@ -469,9 +469,9 @@ def NVVM_ElectSyncOp : NVVM_Op<"elect.sync",
         "{                                  \n"
         ".reg .u32 rx;                      \n"
         ".reg .pred px;                     \n"
-        " mov.u32 %0, 0;                    \n"
+        " mov.pred %0, 0;                   \n"
         "    elect.sync rx | px, 0xFFFFFFFF;\n"
-        "@px mov.u32 %0, 1;                 \n"
+        "@px mov.pred %0, 1;                \n"
         "}\n"
       ); 
     }

mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

@@ -497,9 +497,9 @@ func.func @elect_one_leader_sync() {
   // CHECK: llvm.inline_asm has_side_effects asm_dialect = att "{
   // CHECK-SAME: .reg .u32 rx;
   // CHECK-SAME: .reg .pred px;
-  // CHECK-SAME: mov.u32 $0, 0;
+  // CHECK-SAME: mov.pred $0, 0;
   // CHECK-SAME: elect.sync rx | px, 0xFFFFFFFF;
-  // CHECK-SAME: @px mov.u32 $0, 1;
+  // CHECK-SAME: @px mov.pred $0, 1;
   // CHECK-SAME: "=b"  : () -> i1
   %cnd = nvvm.elect.sync -> i1 
   return