[Mlir-commits] [mlir] [mlir][nvvm] Fix mov.u32 to mov.pred (PR #70027)

[llvmlistbot at llvm.org]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-mlir

Author: Guray Ozen (grypp)

<details>
<summary>Changes</summary>

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`.


---
Full diff: https://github.com/llvm/llvm-project/pull/70027.diff


2 Files Affected:

- (modified) mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td (+2-2) 
- (modified) mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir (+2-2) 


``````````diff
diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
index 9cda7862ccb0fe3..6947cf10e3600d4 100644
--- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
+++ b/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"
       ); 
     }
diff --git a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir b/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
index 3bb0ab90775edf5..b907a86ebc48072 100644
--- a/mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir
+++ b/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 

``````````

</details>


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