diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td index 9cda7862ccb0..6947cf10e360 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 3bb0ab90775e..b907a86ebc48 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