Rename to reductionPredicate

Author: clementval

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

@@ -979,7 +979,7 @@ def NVVM_BarrierOp : NVVM_Op<"barrier", [AttrSizedOperandSegments]> {
 
   let arguments = (ins Optional<I32>:$barrierId, Optional<I32>:$numberOfThreads,
       OptionalAttr<BarrierReductionAttr>:$reductionOp,
-      Optional<I32>:$reductionOperand);
+      Optional<I32>:$reductionPredicate);
   string llvmBuilder = [{
     auto [id, args] = NVVM::BarrierOp::getIntrinsicIDAndArgs(
                         *op, moduleTranslation, builder);
@@ -994,7 +994,7 @@ def NVVM_BarrierOp : NVVM_Op<"barrier", [AttrSizedOperandSegments]> {
 
   let assemblyFormat =
       "(`id` `=` $barrierId^)? (`number_of_threads` `=` $numberOfThreads^)? "
-      "($reductionOp^ $reductionOperand)? (`->` type($res)^)? attr-dict";
+      "($reductionOp^ $reductionPredicate)? (`->` type($res)^)? attr-dict";
 
   let builders = [OpBuilder<(ins), [{
       return build($_builder, $_state, TypeRange{}, Value{}, Value{}, {}, Value{});

mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp

@@ -1518,11 +1518,11 @@ LogicalResult NVVM::BarrierOp::verify() {
     return emitOpError(
         "barrier id is missing, it should be set between 0 to 15");
 
-  if (getBarrierId() && (getReductionOp() || getReductionOperand()))
+  if (getBarrierId() && (getReductionOp() || getReductionPredicate()))
     return emitOpError("reduction are only available when id is 0");
 
-  if ((getReductionOp() && !getReductionOperand()) ||
-      (!getReductionOp() && getReductionOperand()))
+  if ((getReductionOp() && !getReductionPredicate()) ||
+      (!getReductionOp() && getReductionPredicate()))
     return emitOpError("reduction predicate and reduction operation must be "
                        "specified together");
 
@@ -1820,7 +1820,7 @@ mlir::NVVM::IDArgPair NVVM::BarrierOp::getIntrinsicIDAndArgs(
     default:
       llvm_unreachable("Unknown reduction operation for barrier");
     }
-    args.push_back(mt.lookupValue(thisOp.getReductionOperand()));
+    args.push_back(mt.lookupValue(thisOp.getReductionPredicate()));
   } else {
     id = llvm::Intrinsic::nvvm_barrier_cta_sync_aligned_all;
     args.push_back(barrierId);