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

[Guray Ozen]

https://github.com/grypp created https://github.com/llvm/llvm-project/pull/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`.


>From 2660cc30c5f4696c0af7aa9cc3252598cdd463d8 Mon Sep 17 00:00:00 2001
From: Guray Ozen <guray.ozen at gmail.com>
Date: Tue, 24 Oct 2023 12:59:27 +0200
Subject: [PATCH] [mlir][nvvm] Fix mov.u32 to mov.pred

This PR fixes the incorrect PTX instruction. We actually move a predicate not u32, so the correct instruction should be `mov.pred`.
---
 mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td       | 4 ++--
 mlir/test/Conversion/NVVMToLLVM/nvvm-to-llvm.mlir | 4 ++--
 2 files changed, 4 insertions(+), 4 deletions(-)

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