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

[llvmlistbot at llvm.org]

Author: Guray Ozen
Date: 2023-10-25T22:00:09+03:00
New Revision: 01ac180c11d7c16a71ff2de472dc1570dece596c

URL: https://github.com/llvm/llvm-project/commit/01ac180c11d7c16a71ff2de472dc1570dece596c
DIFF: https://github.com/llvm/llvm-project/commit/01ac180c11d7c16a71ff2de472dc1570dece596c.diff

LOG: [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`.

Added: 
    

Modified: 
    mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td
    mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir

Removed: 
    


################################################################################
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