[ODS] Use 'is' instead of 'get' to prefix optional UnitAttr accessors

[amandatang]

Summary

Changing prefixing for optional UnitAttr accessors to follow LLVM convention.

Fixes https://github.com/llvm/llvm-project/issues/54839

[jpienaar]

Nice, overall looks good and reads nicer.

mlir/include/mlir/TableGen/Operator.h

348

Why can't this just be getGetterName and these merged? Computational efficiency?

mlir/lib/Dialect/Linalg/Transforms/ConvertToDestinationStyle.cpp

304

Should this just be isNofold() ? (not sure why attr post fix, not specific to your change though so just OOC)

mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp

1022

Lets end with s/.'/; method will be removed 2023-06-23/

and could send a PR to https://github.com/llvm/mlir-www/blob/main/website/content/deprecation/_index.md to add it there too.

1382

Could you tag all of these with same date so it is easier to find ?

[rriddle]

One thing that's kind of popped up while reading the review is that this going to skew towards attributes with certain naming schemes, which needs to be documented. This is a bit because I feel that 'is' doesn't read well for a number of different cases.

[mehdi_amini]

mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp

291

This does not seem desirable to me that the accessor for getting the attribute name is prefixed with is instead of get.

[amandatang]

mlir/include/mlir/TableGen/Operator.h

348

Yes, the unit attribute check is unnecessary for operands, results and other types. Open to suggestions to improve this logic though.

mlir/lib/Dialect/Linalg/Transforms/ConvertToDestinationStyle.cpp

304

Probably. Not sure where this specific instance of the nofold attribute accessor is defined/generated though.

[amandatang]

In D151821#4385036, @rriddle wrote:

One thing that's kind of popped up while reading the review is that this going to skew towards attributes with certain naming schemes, which needs to be documented. This is a bit because I feel that 'is' doesn't read well for a number of different cases.

Is this just a matter of updating the documentation/style guide or is there a bigger problem regarding naming of unit attributes currently and in the future?

mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp

291

Reverting the changes for the attribute name getters

[amandatang]

Reverted changes for AttrName getters and fixed tests

[aartbik]

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp

828

related to what others said, this reads weird now

hasInserts() would be nice if want to start with a verb

[amandatang]

Use one getGetterName method and logic to handle attributes beginning with has

[amandatang]

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp

828

There are other cases in addition to attribute names that start with "has" where the additional "is" prefix makes it read weird. Handling each case individually does not guarantee consistency and probably isn't feasible. It may be better to update the style guide, so the programmer can ensure proper naming conventions.

[jpienaar]

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp

828

is is a verb. But this goes to River's point below: the prefix works better if one follows naming convention accounting for it.

For boolean accessors is and has are common conventions for prefixing boolean return functions. @lattner commented in the linked issue that in LLVM is is the prefix convention followed.

So 4 options here:

1. Keep the consistent with accessors but inconsistent with manually written code get prefix;
2. Add no prefix and leave it up to user (e.g., just convert to lowerCamel)
3. Enable variable prefix for them ({is, has}, not sure if others are needed ... may ?)
4. Consistently use is as prefix for boolean return functions

For (4) an entry is added to ODS style guide and folks update so that it reads more naturally.

[lattner]

Would it be possible/reasonable to default to "is" and then allow another ODS property to overload it with an arbitrary word like "has"?

[amandatang]

Pull new changes and fix build failure

[jpienaar]

In D151821#4412305, @lattner wrote:

Would it be possible/reasonable to default to "is" and then allow another ODS property to overload it with an arbitrary word like "has"?

It would yes, but there was opposition to enabling accessors to be customized previously (https://github.com/llvm/llvm-project/issues/51258). This is option (3) above though.

It is unclear how much option (3) benefits vs (2) e.g., with (2) you'd have "has_padding" or "is_continuous" in the printed IR rather than "padding" and "continuous", but no need to do anything in ODS or introduce new class for UnitAttr instances to instantiate. And I almost feel that is more readable. We could also do variant of (2) where it is required that there be a prefix in the Attribute ([a-z]*_ or perhaps just an allowed set), else it gets an is.

mlir/docs/DefiningDialects/Operations.md

 the presence of the optional group itself. As such, in these situations the unit
 attribute will not be printed or present in the output and will be automatically
 inferred when parsing by the presence of the optional group itself.
 
 For example, the following operation:
 
 ```tablegen
 def FooOp : ... {
-  let arguments = (ins UnitAttr:$is_read_only);
+  let arguments = (ins UnitAttr:$read_only);
 
-  let assemblyFormat = "attr-dict (`is_read_only` $is_read_only^)?";
+  let assemblyFormat = "attr-dict (`read_only` $read_only^)?";
 }
 ```
 
 would be formatted as such:
 
 ```mlir
 // When the unit attribute is present:
-foo.op is_read_only
+foo.op read_only
 
 // When the unit attribute is not present:
 foo.op
 ```
 
 ##### Optional "else" Group
 
 Optional groups also have support for an "else" group of elements. These are

mlir/include/mlir/Dialect/AMX/AMX.td

let description = [{
     ```mlir
       %0 = amx.tile_muli %a zext, %b zext, %c
         : vector<16x64xi8>, vector<16x64xi8>, vector<16x16xi32>
     ```
   }];
   let arguments = (ins VectorOfRankAndType<[2], [I32, I8]>:$lhs,
                        VectorOfRankAndType<[2], [I32, I8]>:$rhs,
                        VectorOfRankAndType<[2], [I32, I8]>:$acc,
-                       UnitAttr:$isZextLhs,
-                       UnitAttr:$isZextRhs
+                       UnitAttr:$zextLhs,
+                       UnitAttr:$zextRhs
                        );
   let results = (outs VectorOfRankAndType<[2], [I32, I8]>:$res);
   let extraClassDeclaration = [{
     VectorType getLhsVectorType() {
       return ::llvm::cast<VectorType>(getLhs().getType());
     }
     VectorType getRhsVectorType() {
       return ::llvm::cast<VectorType>(getRhs().getType());
     }
     VectorType getVectorType() {
       return ::llvm::cast<VectorType>(getRes().getType());
     }
   }];
-  let assemblyFormat = "$lhs (`zext` $isZextLhs^)? `,` $rhs (`zext` $isZextRhs^)? `,` $acc attr-dict `:` "
+  let assemblyFormat = "$lhs (`zext` $zextLhs^)? `,` $rhs (`zext` $zextRhs^)? `,` $acc attr-dict `:` "
                        "type($lhs) `,` type($rhs) `,` type($acc) ";
   let hasVerifier = 1;
 }
 
 //===----------------------------------------------------------------------===//
 // AMX IntrOp definitions (LLVM compiler facing).
 //===----------------------------------------------------------------------===//
 

mlir/include/mlir/Dialect/EmitC/IR/EmitC.td

let description = [{
 
     Example:
 
     ```mlir
     // Custom form defining the inclusion of `<myheader>`.
     emitc.include <"myheader.h">
 
     // Generic form of the same operation.
-    "emitc.include" (){include = "myheader.h", is_standard_include} : () -> ()
+    "emitc.include" (){include = "myheader.h", standard_include} : () -> ()
 
     // Custom form defining the inclusion of `"myheader"`.
     emitc.include "myheader.h"
 
     // Generic form of the same operation.
     "emitc.include" (){include = "myheader.h"} : () -> ()
     ```
   }];
   let arguments = (ins
     Arg<StrAttr, "source file to include">:$include,
-    UnitAttr:$is_standard_include
+    UnitAttr:$standard_include
   );
   let hasCustomAssemblyFormat = 1;
 }
 
 def EmitC_VariableOp : EmitC_Op<"variable", []> {
   let summary = "Variable operation";
   let description = [{
     The `variable` operation produces an SSA value equal to some value

mlir/include/mlir/Dialect/LLVMIR/LLVMOps.td

let description = [{
     Attempting to define or reference any symbol or any global behavior is
     considered undefined behavior at this time.
   }];
   let arguments = (
     ins Variadic<LLVM_Type>:$operands,
         StrAttr:$asm_string,
         StrAttr:$constraints,
         UnitAttr:$has_side_effects,
-        UnitAttr:$is_align_stack,
+        UnitAttr:$align_stack,
         OptionalAttr<
           DefaultValuedAttr<AsmATTOrIntel, "AsmDialect::AD_ATT">>:$asm_dialect,
         OptionalAttr<ArrayAttr>:$operand_attrs);
 
   let results = (outs Optional<LLVM_Type>:$res);
 
   let assemblyFormat = [{
     (`has_side_effects` $has_side_effects^)?
-    (`is_align_stack` $is_align_stack^)?
+    (`align_stack` $align_stack^)?
     (`asm_dialect` `=` $asm_dialect^)?
     (`operand_attrs` `=` $operand_attrs^)?
     attr-dict
     $asm_string `,` $constraints
     operands `:` functional-type(operands, results)
    }];
 
   let extraClassDeclaration = [{
     static StringRef getElementTypeAttrName() {
       return "elementtype";
     }
   }];
 }
 
 #endif // LLVMIR_OPS

mlir/include/mlir/Dialect/Linalg/TransformOps/LinalgMatchOps.td

let description = !strconcat([{
 
     Succeeds if the specified dimensions satisfy the specified criteria,
     produces a silenceable failure otherwise. Produces a definite failure if
     the operand is not associated with a single payload op.
   }]);
 
   let arguments = (ins TransformHandleTypeInterface:$operand_handle,
                        DenseI64ArrayAttr:$raw_dim_list,
-                       UnitAttr:$is_inverted,
-                       UnitAttr:$is_all,
+                       UnitAttr:$inverted,
+                       UnitAttr:$all,
                        UnitAttr:$parallel,
                        UnitAttr:$reduction);
 
   let results = (outs Optional<TransformParamTypeInterface>:$result);
   let assemblyFormat =
       "$operand_handle `[`"
-      "custom<StructuredTransformDims>($raw_dim_list, $is_inverted, $is_all)"
+      "custom<StructuredTransformDims>($raw_dim_list, $inverted, $all)"
       "`]` attr-dict `:` "
       "custom<SemiFunctionType>(type($operand_handle), type($result))";
 
   let extraClassDeclaration = SingleOpMatcher.extraDeclaration # [{
     ::mlir::DiagnosedSilenceableFailure getDimensionsFor(
         ::mlir::linalg::LinalgOp op,
         ::llvm::SmallVectorImpl<int64_t> &dims);
   }];
class MatchStructuredOperandOp<string opname> : Op<Transform_Dialect, opname, [
     SingleOpMatcher,
     MatchOpInterface,
     MemoryEffectsOpInterface]> {
       
   // TODO: consider an attribute controlling whether to fail or succeed on
   // out-of-bounds accesses.
   let arguments = (ins TransformHandleTypeInterface:$operand_handle,
                        DenseI64ArrayAttr:$raw_position_list,
-                       UnitAttr:$is_inverted,
-                       UnitAttr:$is_all,
+                       UnitAttr:$inverted,
+                       UnitAttr:$all,
                        UnitAttr:$permutation,
                        UnitAttr:$projected_permutation);
 
   // TODO: allow this to bind multiple inputs simultaneously after checking that
   // `transform.foreach` works well in matches.
   let results = (outs Optional<TransformAnyHandle>:$result);
   let assemblyFormat =
       "$operand_handle `[`"
-      "custom<StructuredTransformDims>($raw_position_list, $is_inverted, $is_all)"
+      "custom<StructuredTransformDims>($raw_position_list, $inverted, $all)"
       "`]` attr-dict "
       "`:` custom<SemiFunctionType>(type($operand_handle), type($result))";
 
   let extraClassDeclaration = SingleOpMatcher.extraDeclaration # [{
     ::mlir::DiagnosedSilenceableFailure getPositionsFor(
         ::mlir::linalg::LinalgOp op,
         ::llvm::SmallVectorImpl<int64_t> &positions);
   }];

mlir/include/mlir/Dialect/MLProgram/IR/MLProgramOps.td

let description = [{
     // Mutable global with an undefined initial value.
     ml_program.global mutable @foobar : tensor<?xi32>
     ```
   }];
 
   let arguments = (ins
     SymbolNameAttr:$sym_name,
     TypeAttr:$type,
-    UnitAttr:$is_mutable,
+    UnitAttr:$global_mutable,
     OptionalAttr<AnyAttr>:$value,
     OptionalAttr<StrAttr>:$sym_visibility
   );
 
   let assemblyFormat = [{
     custom<SymbolVisibility>($sym_visibility)
-    (`mutable` $is_mutable^)?
+    (`mutable` $global_mutable^)?
     $sym_name ``
     custom<TypedInitialValue>($type, $value)
     attr-dict
   }];
 
   let hasVerifier = 1;
 }
 

mlir/include/mlir/TableGen/Operator.h

public:
   OperandOrAttribute getArgToOperandOrAttribute(int index) const;
 
   /// Returns the builders of this operation.
   ArrayRef<Builder> getBuilders() const { return builders; }
 
   /// Returns the getter name for the accessor of `name`.
   std::string getGetterName(StringRef name) const;
 
+  /// Returns the getter name for the accessor of `name`.
+  /// Searches for the `name` in attributes to check for the UnitAttr type.
+  std::string getGetterNameUnitAttr(StringRef name) const;
+
+  /// Returns the getter name for the accessor of `name`.
+  std::string getGetterNameAttr(const NamedAttribute &attr) const;
+
   /// Returns the setter name for the accessor of `name`.
   std::string getSetterName(StringRef name) const;
 
   /// Returns the remove name for the accessor of `name`.
   std::string getRemoverName(StringRef name) const;
 
   bool hasFolder() const;
 

mlir/lib/Conversion/AMDGPUToROCDL/AMDGPUToROCDL.cpp

static std::optional<StringRef> mfmaOpToIntrinsic(MFMAOp mfma,
   Type sourceElem = mfma.getSourceA().getType();
   if (auto sourceType = dyn_cast<VectorType>(sourceElem))
     sourceElem = sourceType.getElementType();
   Type destElem = mfma.getDestC().getType();
   if (auto destType = dyn_cast<VectorType>(destElem))
     destElem = destType.getElementType();
 
   if (sourceElem.isF32() && destElem.isF32()) {
-    if (mfma.getReducePrecision() && chipset.minorVersion >= 0x40) {
+    if (mfma.isReducePrecision() && chipset.minorVersion >= 0x40) {
       if (m == 32 && n == 32 && k == 4 && b == 1)
         return ROCDL::mfma_f32_32x32x4_xf32::getOperationName();
       if (m == 16 && n == 16 && k == 8 && b == 1)
         return ROCDL::mfma_f32_16x16x8_xf32::getOperationName();
     }
     if (m == 32 && n == 32 && k == 1 && b == 2)
       return ROCDL::mfma_f32_32x32x1f32::getOperationName();
     if (m == 16 && n == 16 && k == 1 && b == 4)
struct MFMAOpLowering : public ConvertOpToLLVMPattern<MFMAOp> {
   matchAndRewrite(MFMAOp op, MFMAOpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Location loc = op.getLoc();
     Type outType = typeConverter->convertType(op.getDestD().getType());
 
     if (chipset.majorVersion != 9 || chipset.minorVersion < 0x08)
       return op->emitOpError("MFMA only supported on gfx908+");
     uint32_t getBlgpField = static_cast<uint32_t>(op.getBlgp());
-    if (op.getNegateA() || op.getNegateB() || op.getNegateC()) {
+    if (op.isNegateA() || op.isNegateB() || op.isNegateC()) {
       if (chipset.minorVersion < 0x40)
         return op.emitOpError("negation unsupported on older than gfx840");
       getBlgpField |=
-          op.getNegateA() | (op.getNegateB() << 1) | (op.getNegateC() << 2);
+          op.isNegateA() | (op.isNegateB() << 1) | (op.isNegateC() << 2);
     }
     std::optional<StringRef> maybeIntrinsic = mfmaOpToIntrinsic(op, chipset);
     if (!maybeIntrinsic.has_value())
       return op.emitOpError("no intrinsic matching MFMA size on given chipset");
     OperationState loweredOp(loc, *maybeIntrinsic);
     loweredOp.addTypes(outType);
     loweredOp.addOperands(
         {mfmaConcatIfNeeded(rewriter, loc, adaptor.getSourceA()),

mlir/lib/Conversion/GPUCommon/GPUToLLVMConversion.cpp

LogicalResult ConvertHostUnregisterOpToGpuRuntimeCallPattern::matchAndRewrite(
 
   rewriter.eraseOp(op);
   return success();
 }
 
 LogicalResult ConvertAllocOpToGpuRuntimeCallPattern::matchAndRewrite(
     gpu::AllocOp allocOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
-  if (adaptor.getHostShared())
+  if (adaptor.isHostShared())
     return rewriter.notifyMatchFailure(
         allocOp, "host_shared allocation is not supported");
 
   MemRefType memRefType = allocOp.getType();
 
   if (failed(areAllLLVMTypes(allocOp, adaptor.getOperands(), rewriter)) ||
       !isConvertibleAndHasIdentityMaps(memRefType) ||
       failed(isAsyncWithOneDependency(rewriter, allocOp)))

mlir/lib/Conversion/GPUToNVVM/LowerGpuOpsToNVVMOps.cpp

 /// must be run by the entire subgroup, otherwise it is undefined behaviour.
 struct GPUSubgroupReduceOpLowering
     : public ConvertOpToLLVMPattern<gpu::SubgroupReduceOp> {
   using ConvertOpToLLVMPattern<gpu::SubgroupReduceOp>::ConvertOpToLLVMPattern;
   LogicalResult
 
   matchAndRewrite(gpu::SubgroupReduceOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    if (!op.getUniform())
+    if (!op.isUniform())
       return rewriter.notifyMatchFailure(
           op, "cannot be lowered to redux as the op must be run "
               "uniformly (entire subgroup).");
     if (!op.getValue().getType().isInteger(32))
       return rewriter.notifyMatchFailure(op, "unsupported data type");
 
     std::optional<NVVM::ReduxKind> mode = convertReduxKind(op.getOp());
     if (!mode.has_value())

mlir/lib/Conversion/GPUToNVVM/WmmaOpsToNvvm.cpp

matchAndRewrite(gpu::SubgroupMmaLoadMatrixOp subgroupMmaLoadMatrixOp,
                   OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Operation *op = subgroupMmaLoadMatrixOp.getOperation();
     if (failed(areAllLLVMTypes(op, adaptor.getOperands(), rewriter)))
       return failure();
 
     // Get the shape of the MMAMatrix type being returned. The shape will
     // choose which intrinsic this op will be lowered to.
-    NVVM::MMALayout layout = subgroupMmaLoadMatrixOp.getTranspose()
+    NVVM::MMALayout layout = subgroupMmaLoadMatrixOp.isTranspose()
                                  ? NVVM::MMALayout::col
                                  : NVVM::MMALayout::row;
     gpu::MMAMatrixType retType =
         cast<gpu::MMAMatrixType>(subgroupMmaLoadMatrixOp.getRes().getType());
     ArrayRef<int64_t> retTypeShape = retType.getShape();
     int64_t m = 0;
     int64_t n = 0;
     int64_t k = 0;
matchAndRewrite(gpu::SubgroupMmaStoreMatrixOp subgroupMmaStoreMatrixOp,
     Location loc = op->getLoc();
 
     SmallVector<Value, 4> storeOpOperands;
     // Get the shape of the MMAMatrix type being stored. The shape will
     // choose which intrinsic this op will be lowered to.
     gpu::MMAMatrixType srcType =
         cast<gpu::MMAMatrixType>(subgroupMmaStoreMatrixOp.getSrc().getType());
     ArrayRef<int64_t> srcTypeShape = srcType.getShape();
-    NVVM::MMALayout layout = subgroupMmaStoreMatrixOp.getTranspose()
+    NVVM::MMALayout layout = subgroupMmaStoreMatrixOp.isTranspose()
                                  ? NVVM::MMALayout::col
                                  : NVVM::MMALayout::row;
     NVVM::MMATypes eltype = getElementType(srcType);
     int64_t m = srcTypeShape[0];
     int64_t n = srcTypeShape[1];
     int64_t k = NVVM::WMMAStoreOp::inferKDimension(m, n, eltype);
     if (NVVM::WMMAStoreOp::getIntrinsicID(m, n, k, layout, eltype) == 0)
       return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
gpu::MMAMatrixType aType =
         cast<gpu::MMAMatrixType>(subgroupMmaComputeOp.getOpA().getType());
     ArrayRef<int64_t> aTypeShape = aType.getShape();
     gpu::MMAMatrixType cType =
         cast<gpu::MMAMatrixType>(subgroupMmaComputeOp.getOpC().getType());
     ArrayRef<int64_t> cTypeShape = cType.getShape();
     int64_t m = cTypeShape[0];
     int64_t n = cTypeShape[1];
     int64_t k = aTypeShape[1];
-    NVVM::MMALayout aLayout = subgroupMmaComputeOp.getATranspose()
+    NVVM::MMALayout aLayout = subgroupMmaComputeOp.isATranspose()
                                   ? NVVM::MMALayout::col
                                   : NVVM::MMALayout::row;
-    NVVM::MMALayout bLayout = subgroupMmaComputeOp.getBTranspose()
+    NVVM::MMALayout bLayout = subgroupMmaComputeOp.isBTranspose()
                                   ? NVVM::MMALayout::col
                                   : NVVM::MMALayout::row;
     NVVM::MMATypes sourceType = getElementType(aType);
     NVVM::MMATypes destType = getElementType(cType);
     if (NVVM::WMMAMmaOp::getIntrinsicID(m, n, k, aLayout, bLayout, sourceType,
                                         destType) == 0)
       return rewriter.notifyMatchFailure(op, kInvalidCaseStr);
 

mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp

matchAndRewrite(gpu::AllReduceOp op, OpAdaptor adaptor,
 
     // gpu.all_reduce can have either reduction op attribute or reduction
     // region. Only attribute version is supported.
     if (!opType)
       return failure();
 
     auto result =
         createGroupReduceOp(rewriter, op.getLoc(), adaptor.getValue(), *opType,
-                            /*isGroup*/ true, op.getUniform());
+                            /*isGroup*/ true, op.isUniform());
     if (!result)
       return failure();
 
     rewriter.replaceOp(op, *result);
     return success();
   }
 };
 
 /// Pattern to convert a gpu.subgroup_reduce op into a SPIR-V group op.
 class GPUSubgroupReduceConversion final
     : public OpConversionPattern<gpu::SubgroupReduceOp> {
 public:
   using OpConversionPattern::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(gpu::SubgroupReduceOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     auto opType = op.getOp();
     auto result =
         createGroupReduceOp(rewriter, op.getLoc(), adaptor.getValue(), opType,
-                            /*isGroup*/ false, op.getUniform());
+                            /*isGroup*/ false, op.isUniform());
     if (!result)
       return failure();
 
     rewriter.replaceOp(op, *result);
     return success();
   }
 };
 

mlir/lib/Conversion/GPUToSPIRV/WmmaOpsToSPIRV.cpp

matchAndRewrite(gpu::SubgroupMmaLoadMatrixOp subgroupMmaLoadMatrixOp,
     Value bufferPtr = spirv::getElementPtr(
         *getTypeConverter<SPIRVTypeConverter>(), memrefType,
         adaptor.getSrcMemref(), adaptor.getIndices(), loc, rewriter);
     auto coopType = convertMMAToSPIRVType(retType);
     int64_t stride = subgroupMmaLoadMatrixOp.getLeadDimension().getSExtValue();
     auto i32Type = rewriter.getI32Type();
     auto strideValue = rewriter.create<spirv::ConstantOp>(
         loc, i32Type, IntegerAttr::get(i32Type, stride));
-    bool isColMajor = static_cast<bool>(subgroupMmaLoadMatrixOp.getTranspose());
+    bool isColMajor = static_cast<bool>(subgroupMmaLoadMatrixOp.isTranspose());
     auto columnMajor = rewriter.create<spirv::ConstantOp>(
         loc, rewriter.getI1Type(), rewriter.getBoolAttr(isColMajor));
     rewriter.replaceOpWithNewOp<spirv::NVCooperativeMatrixLoadOp>(
         subgroupMmaLoadMatrixOp, coopType, bufferPtr, strideValue, columnMajor,
         spirv::MemoryAccessAttr());
     return success();
   }
 };
matchAndRewrite(gpu::SubgroupMmaStoreMatrixOp subgroupMmaStoreMatrixOp,
     Value bufferPtr = spirv::getElementPtr(
         *getTypeConverter<SPIRVTypeConverter>(), memrefType,
         adaptor.getDstMemref(), adaptor.getIndices(), loc, rewriter);
     int64_t stride = subgroupMmaStoreMatrixOp.getLeadDimension().getSExtValue();
     auto i32Type = rewriter.getI32Type();
     auto strideValue = rewriter.create<spirv::ConstantOp>(
         loc, i32Type, IntegerAttr::get(i32Type, stride));
     bool useColMajor =
-        static_cast<bool>(subgroupMmaStoreMatrixOp.getTranspose());
+        static_cast<bool>(subgroupMmaStoreMatrixOp.isTranspose());
     auto columnMajor = rewriter.create<spirv::ConstantOp>(
         loc, rewriter.getI1Type(), rewriter.getBoolAttr(useColMajor));
     rewriter.replaceOpWithNewOp<spirv::NVCooperativeMatrixStoreOp>(
         subgroupMmaStoreMatrixOp, bufferPtr, adaptor.getSrc(), strideValue,
         columnMajor, spirv::MemoryAccessAttr());
     return success();
   }
 };

mlir/lib/Conversion/MemRefToLLVM/MemRefToLLVM.cpp

matchAndRewrite(memref::GlobalOp global, OpAdaptor adaptor,
 
     uint64_t alignment = global.getAlignment().value_or(0);
     FailureOr<unsigned> addressSpace =
         getTypeConverter()->getMemRefAddressSpace(type);
     if (failed(addressSpace))
       return global.emitOpError(
           "memory space cannot be converted to an integer address space");
     auto newGlobal = rewriter.replaceOpWithNewOp<LLVM::GlobalOp>(
-        global, arrayTy, global.getConstant(), linkage, global.getSymName(),
+        global, arrayTy, global.isConstant(), linkage, global.getSymName(),
         initialValue, alignment, *addressSpace);
     if (!global.isExternal() && global.isUninitialized()) {
       Block *blk = new Block();
       newGlobal.getInitializerRegion().push_back(blk);
       rewriter.setInsertionPointToStart(blk);
       Value undef[] = {
           rewriter.create<LLVM::UndefOp>(global.getLoc(), arrayTy)};
       rewriter.create<LLVM::ReturnOp>(global.getLoc(), undef);

mlir/lib/Conversion/NVGPUToNVVM/NVGPUToNVVM.cpp

matchAndRewrite(nvgpu::MmaSyncOp op, OpAdaptor adaptor,
     // choose which intrinsic this op will be lowered to.
     VectorType aType = op.getMatrixA().getType();
     VectorType bType = op.getMatrixA().getType();
     VectorType cType = op.getMatrixC().getType();
 
     std::array<int64_t, 3> gemmShape = op.getMmaShapeAsArray();
 
     // Tensor Cores (mma.sync) on F32 works only with TensorFloat32 (TF32).
     bool tf32Enabled = op->hasAttr(op.getTf32EnabledAttrName());
     if (aType.getElementType().isF32() && !tf32Enabled)
       return failure();
 
     FailureOr<NVVM::MMATypes> ptxTypeA = getNvvmMmaType(aType);
     if (failed(ptxTypeA))
       return op->emitOpError("failed to deduce operand PTX types");
     FailureOr<NVVM::MMATypes> ptxTypeB = getNvvmMmaType(bType);
     if (failed(ptxTypeB))
matchAndRewrite(nvgpu::DeviceAsyncCopyOp op, OpAdaptor adaptor,
     scrPtr = rewriter.create<LLVM::AddrSpaceCastOp>(loc, srcPointerGlobalType,
                                                     scrPtr);
     int64_t dstElements = adaptor.getDstElements().getZExtValue();
     int64_t sizeInBytes =
         (dstMemrefType.getElementTypeBitWidth() * dstElements) / 8;
     // bypass L1 is only supported for byte sizes of 16, we drop the hint
     // otherwise.
     UnitAttr bypassL1 =
-        sizeInBytes == 16 ? adaptor.getBypassL1Attr() : UnitAttr();
+        sizeInBytes == 16 ? adaptor.isBypassL1Attr() : UnitAttr();
 
     // When the optional SrcElements argument is present, the source (global
     // memory) of CpAsyncOp is read only for SrcElements number of elements. The
     // rest of the DstElements in the destination (shared memory) are filled
     // with zeros.
     if (op.getSrcElements())
       emitCpAsyncOpZfillAsm(loc, dstPtr, scrPtr,
                             rewriter.create<LLVM::ConstantOp>(

mlir/lib/Dialect/AMDGPU/IR/AMDGPUDialect.cpp

return emitOpError(
         "double-precision ops do not support permuting lanes of B");
   if (destElem.isF64() && getCbsz() != 0)
     return emitOpError(
         "double-precision ops do not support permuting lanes of A");
   if (getAbid() >= (1u << getCbsz()))
     return emitOpError(
         "block ID for permuting A (abid) must be below 2 ** cbsz");
 
-  if ((getNegateA() || getNegateB() || getNegateC()) && !destElem.isF64())
+  if ((isNegateA() || isNegateB() || isNegateC()) && !destElem.isF64())
     return emitOpError(
         "negation flags only available for double-precision operations");
 
   return success();
 }
 
 #include "mlir/Dialect/AMDGPU/IR/AMDGPUEnums.cpp.inc"
 
 #define GET_ATTRDEF_CLASSES
 #include "mlir/Dialect/AMDGPU/IR/AMDGPUAttributes.cpp.inc"
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/AMDGPU/IR/AMDGPU.cpp.inc"

mlir/lib/Dialect/AMX/Transforms/LegalizeForLLVMExport.cpp

matchAndRewrite(TileMulIOp op, OpAdaptor adaptor,
     VectorType cType = op.getVectorType();
     // Determine m x n x k tile sizes.
     std::pair<Value, Value> tsza =
         getTileSizes(rewriter, *getTypeConverter(), aType, op.getLoc());
     std::pair<Value, Value> tszb =
         getTileSizes(rewriter, *getTypeConverter(), bType, op.getLoc());
     // Replace operation with intrinsic.
     Type resType = typeConverter->convertType(cType);
-    bool zexta = op.getIsZextLhs();
-    bool zextb = op.getIsZextRhs();
+    bool zexta = op.isZextLhs();
+    bool zextb = op.isZextRhs();
     if (zexta && zextb)
       rewriter.replaceOpWithNewOp<amx::x86_amx_tdpbuud>(
           op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(),
           adaptor.getLhs(), adaptor.getRhs());
     else if (zexta && !zextb)
       rewriter.replaceOpWithNewOp<amx::x86_amx_tdpbusd>(
           op, resType, tsza.first, tszb.second, tsza.second, adaptor.getAcc(),
           adaptor.getLhs(), adaptor.getRhs());

mlir/lib/Dialect/Bufferization/IR/BufferizationOps.cpp

LogicalResult DeallocTensorOp::bufferize(RewriterBase &rewriter,
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // ToTensorOp
 //===----------------------------------------------------------------------===//
 
 bool ToTensorOp::isWritable(Value value, const AnalysisState &state) {
-  return getWritable();
+  return isWritable();
 }
 
 OpFoldResult ToTensorOp::fold(FoldAdaptor) {
   if (auto toMemref = getMemref().getDefiningOp<ToMemrefOp>())
     // Approximate alias analysis by conservatively folding only when no there
     // is no interleaved operation.
     if (toMemref->getBlock() == this->getOperation()->getBlock() &&
         toMemref->getNextNode() == this->getOperation())

mlir/lib/Dialect/Bufferization/Transforms/OneShotAnalysis.cpp

if (isa<ToMemrefOp>(op.getOperation())) {
       op->emitError("to_memref ops are not supported by One-Shot Analysis");
       return WalkResult::interrupt();
     }
 
     // Input IR may not contain any ToTensorOps without the "restrict"
     // attribute. Such tensors may alias any other tensor, which is currently
     // not handled in the analysis.
     if (auto toTensorOp = dyn_cast<ToTensorOp>(op.getOperation())) {
-      if (!toTensorOp.getRestrict()) {
+      if (!toTensorOp.isRestrict()) {
         op->emitError("to_tensor ops without `restrict` are not supported by "
                       "One-Shot Analysis");
         return WalkResult::interrupt();
       }
     }
 
     for (OpOperand &opOperand : op->getOpOperands()) {
       if (isa<TensorType>(opOperand.get().getType())) {

mlir/lib/Dialect/EmitC/IR/EmitC.cpp

OpFoldResult emitc::ConstantOp::fold(FoldAdaptor adaptor) {
   return getValue();
 }
 
 //===----------------------------------------------------------------------===//
 // IncludeOp
 //===----------------------------------------------------------------------===//
 
 void IncludeOp::print(OpAsmPrinter &p) {
-  bool standardInclude = getIsStandardInclude();
+  bool standardInclude = isStandardInclude();
 
   p << " ";
   if (standardInclude)
     p << "<";
   p << "\"" << getInclude() << "\"";
   if (standardInclude)
     p << ">";
 }
 
 ParseResult IncludeOp::parse(OpAsmParser &parser, OperationState &result) {
   bool standardInclude = !parser.parseOptionalLess();
 
   StringAttr include;
   OptionalParseResult includeParseResult =
       parser.parseOptionalAttribute(include, "include", result.attributes);
   if (!includeParseResult.has_value())
     return parser.emitError(parser.getNameLoc()) << "expected string attribute";
 
   if (standardInclude && parser.parseOptionalGreater())
     return parser.emitError(parser.getNameLoc())
            << "expected trailing '>' for standard include";
 
   if (standardInclude)
-    result.addAttribute("is_standard_include",
-                        UnitAttr::get(parser.getContext()));
+    result.addAttribute("standard_include", UnitAttr::get(parser.getContext()));
 
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // VariableOp
 //===----------------------------------------------------------------------===//
 

mlir/lib/Dialect/GPU/IR/GPUDialect.cpp

static bool canMakeGroupOpUniform(Operation *op) {
   Region &body = launchOp.getBody();
   assert(!body.empty() && "Invalid region");
 
   // Only convert ops in gpu::launch entry block for now.
   return op->getBlock() == &body.front();
 }
 
 OpFoldResult gpu::AllReduceOp::fold(FoldAdaptor /*adaptor*/) {
-  if (!getUniform() && canMakeGroupOpUniform(*this)) {
+  if (!isUniform() && canMakeGroupOpUniform(*this)) {
     setUniform(true);
     return getResult();
   }
 
   return nullptr;
 }
 
 // TODO: Support optional custom attributes (without dialect prefix).
LogicalResult gpu::SubgroupReduceOp::verify() {
   if (!verifyReduceOpAndType(opName, getType())) {
     return emitError() << '`' << gpu::stringifyAllReduceOperation(opName)
                        << "` accumulator is only compatible with Integer type";
   }
   return success();
 }
 
 OpFoldResult gpu::SubgroupReduceOp::fold(FoldAdaptor /*adaptor*/) {
-  if (!getUniform() && canMakeGroupOpUniform(*this)) {
+  if (!isUniform() && canMakeGroupOpUniform(*this)) {
     setUniform(true);
     return getResult();
   }
 
   return nullptr;
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/GPU/TransformOps/GPUTransformOps.cpp

 DiagnosedSilenceableFailure
 transform::MapForallToBlocks::applyToOne(Operation *target,
                                          ApplyToEachResultList &results,
                                          transform::TransformState &state) {
   LaunchOp gpuLaunch = dyn_cast<LaunchOp>(target);
   IRRewriter rewriter(getContext());
   auto transformOp = cast<TransformOpInterface>(getOperation());
 
-  if (!getGenerateGpuLaunch() && !gpuLaunch) {
+  if (!isGenerateGpuLaunch() && !gpuLaunch) {
     DiagnosedSilenceableFailure diag =
         emitSilenceableError()
         << "Given target is not gpu.launch, set `generate_gpu_launch` "
            "attribute";
     diag.attachNote(target->getLoc()) << "when applied to this payload op";
     return diag;
   }
 
   scf::ForallOp topLevelForallOp;
   DiagnosedSilenceableFailure diag = mlir::transform::gpu::findTopLevelForallOp(
       target, topLevelForallOp, transformOp);
   if (!diag.succeeded()) {
     diag.attachNote(target->getLoc()) << "when applied to this payload op";
     return diag;
   }
 
   SmallVector<int64_t> gridDims{getGridDims()};
-  if (!getGenerateGpuLaunch() && gridDims.size() != 3)
+  if (!isGenerateGpuLaunch() && gridDims.size() != 3)
     return transformOp.emitDefiniteFailure("transform require size-3 mapping");
 
   OpBuilder::InsertionGuard guard(rewriter);
   rewriter.setInsertionPoint(topLevelForallOp);
 
   // Generate gpu launch here and move the forall inside
-  if (getGenerateGpuLaunch()) {
+  if (isGenerateGpuLaunch()) {
     DiagnosedSilenceableFailure diag =
         createGpuLaunch(rewriter, target->getLoc(), transformOp, gpuLaunch);
     if (!diag.succeeded()) {
       return diag;
     }
     rewriter.setInsertionPointToStart(&gpuLaunch.getBody().front());
     Operation *newForallOp = rewriter.clone(*topLevelForallOp);
     rewriter.eraseOp(topLevelForallOp);

mlir/lib/Dialect/GPU/Transforms/AllReduceLowering.cpp

explicit GpuAllReduceConversion(MLIRContext *context)
       : RewritePattern(gpu::GPUFuncOp::getOperationName(), 1, context) {}
 
   LogicalResult matchAndRewrite(Operation *op,
                                 PatternRewriter &rewriter) const override {
     auto funcOp = cast<gpu::GPUFuncOp>(op);
 
     SmallVector<gpu::AllReduceOp> reduceOps;
     auto callback = [&](gpu::AllReduceOp reduceOp) -> WalkResult {
-      if (!reduceOp.getUniform())
+      if (!reduceOp.isUniform())
         return WalkResult::interrupt();
 
       reduceOps.emplace_back(reduceOp);
       return WalkResult::advance();
     };
 
     if (funcOp.walk(callback).wasInterrupted() || reduceOps.empty())
       return rewriter.notifyMatchFailure(

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

 void AllocaOp::print(OpAsmPrinter &p) {
   Type elemTy = llvm::cast<LLVM::LLVMPointerType>(getType()).getElementType();
   if (!elemTy)
     elemTy = *getElemType();
 
   auto funcTy =
       FunctionType::get(getContext(), {getArraySize().getType()}, {getType()});
 
-  if (getInalloca())
+  if (isInalloca())
     p << " inalloca";
 
   p << ' ' << getArraySize() << " x " << elemTy;
   if (getAlignment() && *getAlignment() != 0)
     p.printOptionalAttrDict((*this)->getAttrs(),
                             {kElemTypeAttrName, getInallocaAttrName()});
   else
     p.printOptionalAttrDict(
void GEPOp::build(OpBuilder &builder, OperationState &result, Type resultType,
   SmallVector<Value> dynamicIndices;
   destructureIndices(elementType, indices, rawConstantIndices, dynamicIndices);
 
   result.addTypes(resultType);
   result.addAttributes(attributes);
   result.addAttribute(getRawConstantIndicesAttrName(result.name),
                       builder.getDenseI32ArrayAttr(rawConstantIndices));
   if (inbounds) {
-    result.addAttribute(getInboundsAttrName(result.name),
-                        builder.getUnitAttr());
+    result.addAttribute(getInboundsAttrName(result.name), builder.getUnitAttr());
   }
   if (llvm::cast<LLVMPointerType>(extractVectorElementType(basePtr.getType()))
           .isOpaque())
     result.addAttribute(kElemTypeAttrName, TypeAttr::get(elementType));
   result.addOperands(basePtr);
   result.addOperands(dynamicIndices);
 }
 
if (LLVMFuncOp func = (*this)->getParentOfType<LLVMFuncOp>()) {
     if (!func.getPersonality())
       return emitError(
           "llvm.landingpad needs to be in a function with a personality");
   }
 
   // Consistency of llvm.landingpad result types is checked in
   // LLVMFuncOp::verify().
 
-  if (!getCleanup() && getOperands().empty())
+  if (!isCleanup() && getOperands().empty())
     return emitError("landingpad instruction expects at least one clause or "
                      "cleanup attribute");
 
   for (unsigned idx = 0, ie = getNumOperands(); idx < ie; idx++) {
     value = getOperand(idx);
     bool isFilter = llvm::isa<LLVMArrayType>(value.getType());
     if (isFilter) {
       // FIXME: Verify filter clauses when arrays are appropriately handled
if (isFilter) {
       return emitError("clause #")
              << idx << " is not a known constant - null, addressof, bitcast";
     }
   }
   return success();
 }
 
 void LandingpadOp::print(OpAsmPrinter &p) {
-  p << (getCleanup() ? " cleanup " : " ");
+  p << (isCleanup() ? " cleanup " : " ");
 
   // Clauses
   for (auto value : getOperands()) {
     // Similar to llvm - if clause is an array type then it is filter
     // clause else catch clause
     bool isArrayTy = llvm::isa<LLVMArrayType>(value.getType());
     p << '(' << (isArrayTy ? "filter " : "catch ") << value << " : "
       << value.getType() << ") ";
void GlobalOp::build(OpBuilder &builder, OperationState &result, Type type,
                      bool isConstant, Linkage linkage, StringRef name,
                      Attribute value, uint64_t alignment, unsigned addrSpace,
                      bool dsoLocal, bool threadLocal,
                      ArrayRef<NamedAttribute> attrs) {
   result.addAttribute(getSymNameAttrName(result.name),
                       builder.getStringAttr(name));
   result.addAttribute(getGlobalTypeAttrName(result.name), TypeAttr::get(type));
   if (isConstant)
-    result.addAttribute(getConstantAttrName(result.name),
-                        builder.getUnitAttr());
+    result.addAttribute(getConstantAttrName(result.name), builder.getUnitAttr());
   if (value)
     result.addAttribute(getValueAttrName(result.name), value);
   if (dsoLocal)
-    result.addAttribute(getDsoLocalAttrName(result.name),
-                        builder.getUnitAttr());
+    result.addAttribute(getDsoLocalAttrName(result.name), builder.getUnitAttr());
   if (threadLocal)
     result.addAttribute(getThreadLocal_AttrName(result.name),
                         builder.getUnitAttr());
 
   // Only add an alignment attribute if the "alignment" input
   // is different from 0. The value must also be a power of two, but
   // this is tested in GlobalOp::verify, not here.
   if (alignment != 0)
void GlobalOp::build(OpBuilder &builder, OperationState &result, Type type,
   result.addRegion();
 }
 
 void GlobalOp::print(OpAsmPrinter &p) {
   p << ' ' << stringifyLinkage(getLinkage()) << ' ';
   StringRef visibility = stringifyVisibility(getVisibility_());
   if (!visibility.empty())
     p << visibility << ' ';
-  if (getThreadLocal_())
+  if (isThreadLocal_())
     p << "thread_local ";
   if (auto unnamedAddr = getUnnamedAddr()) {
     StringRef str = stringifyUnnamedAddr(*unnamedAddr);
     if (!str.empty())
       p << str << ' ';
   }
-  if (getConstant())
+  if (isConstant())
     p << "constant ";
   p.printSymbolName(getSymName());
   p << '(';
   if (auto value = getValueOrNull())
     p.printAttribute(value);
   p << ')';
   // Note that the alignment attribute is printed using the
   // default syntax here, even though it is an inherent attribute
void LLVMFuncOp::build(OpBuilder &builder, OperationState &result,
   result.addAttribute(getFunctionTypeAttrName(result.name),
                       TypeAttr::get(type));
   result.addAttribute(getLinkageAttrName(result.name),
                       LinkageAttr::get(builder.getContext(), linkage));
   result.addAttribute(getCConvAttrName(result.name),
                       CConvAttr::get(builder.getContext(), cconv));
   result.attributes.append(attrs.begin(), attrs.end());
   if (dsoLocal)
-    result.addAttribute(getDsoLocalAttrName(result.name),
-                        builder.getUnitAttr());
+    result.addAttribute(getDsoLocalAttrName(result.name), builder.getUnitAttr());
   if (functionEntryCount)
     result.addAttribute(getFunctionEntryCountAttrName(result.name),
                         builder.getI64IntegerAttr(functionEntryCount.value()));
   if (argAttrs.empty())
     return;
 
   assert(llvm::cast<LLVMFunctionType>(type).getNumParams() == argAttrs.size() &&
          "expected as many argument attribute lists as arguments");

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

bool LLVM::LoadOp::canUsesBeRemoved(
   if (blockingUses.size() != 1)
     return false;
   Value blockingUse = (*blockingUses.begin())->get();
   // If the blocking use is the slot ptr itself, there will be enough
   // context to reconstruct the result of the load at removal time, so it can
   // be removed (provided it loads the exact stored value and is not
   // volatile).
   return blockingUse == slot.ptr && getAddr() == slot.ptr &&
-         getResult().getType() == slot.elemType && !getVolatile_();
+         getResult().getType() == slot.elemType && !isVolatile_();
 }
 
 DeletionKind LLVM::LoadOp::removeBlockingUses(
     const MemorySlot &slot, const SmallPtrSetImpl<OpOperand *> &blockingUses,
     RewriterBase &rewriter, Value reachingDefinition) {
   // `canUsesBeRemoved` checked this blocking use must be the loaded slot
   // pointer.
   rewriter.replaceAllUsesWith(getResult(), reachingDefinition);
   return DeletionKind::Delete;
 }
 
 bool LLVM::StoreOp::canUsesBeRemoved(
     const MemorySlot &slot, const SmallPtrSetImpl<OpOperand *> &blockingUses,
     SmallVectorImpl<OpOperand *> &newBlockingUses) {
   if (blockingUses.size() != 1)
     return false;
   Value blockingUse = (*blockingUses.begin())->get();
   // If the blocking use is the slot ptr itself, dropping the store is
   // fine, provided we are currently promoting its target value. Don't allow a
   // store OF the slot pointer, only INTO the slot pointer.
   return blockingUse == slot.ptr && getAddr() == slot.ptr &&
          getValue() != slot.ptr && getValue().getType() == slot.elemType &&
-         !getVolatile_();
+         !isVolatile_();
 }
 
 DeletionKind LLVM::StoreOp::removeBlockingUses(
     const MemorySlot &slot, const SmallPtrSetImpl<OpOperand *> &blockingUses,
     RewriterBase &rewriter, Value reachingDefinition) {
   // `canUsesBeRemoved` checked this blocking use must be the stored slot
   // pointer.
   for (Operation *user : slot.ptr.getUsers())

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

return failure(parser.parseOperandList(ops) ||
                                         parser.getNameLoc(), result.operands));
 }
 
 void VoteBallotOp::print(OpAsmPrinter &p) { printNVVMIntrinsicOp(p, *this); }
 
 LogicalResult CpAsyncOp::verify() {
   if (getSize() != 4 && getSize() != 8 && getSize() != 16)
     return emitError("expected byte size to be either 4, 8 or 16.");
-  if (getBypassL1() && getSize() != 16)
+  if (isBypassL1() && getSize() != 16)
     return emitError("bypass l1 is only support for 16 bytes copy.");
   return success();
 }
 
 // Given the element type of an operand and whether or not it is an accumulator,
 // this function returns the PTX type (`NVVM::MMATypes`) that corresponds to the
 // operand's element type.
 std::optional<mlir::NVVM::MMATypes>

mlir/lib/Dialect/Linalg/TransformOps/LinalgMatchOps.cpp

if (!matchReduction(linalgOp.getRegionOutputArgs(), *position,
                         combinerOps)) {
       return emitSilenceableError() << "could not match reduction";
     }
     if (combinerOps.size() != 1) {
       return emitSilenceableError() << "reduction combiner is not a single op";
     }
     return DiagnosedSilenceableFailure::success();
   }
-  if (getPassthrough()) {
+  if (isPassthrough()) {
     Block &body = linalgOp->getRegion(0).front();
     if (body.getTerminator()->getOperands() != linalgOp.getRegionInputArgs()) {
       return emitSilenceableError() << "not a passthrough";
     }
     return DiagnosedSilenceableFailure::success();
   }
   return emitDefiniteFailure() << "unknown body condition";
 }
 
 LogicalResult transform::MatchStructuredBodyOp::verify() {
-  if (getReductionPosition() && getPassthrough()) {
+  if (getReductionPosition() && isPassthrough()) {
     return emitOpError() << "reduction position and passthrough conditions are "
                             "mutually exclusive";
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // Utilities for structured match predicates.
DiagnosedSilenceableFailure transform::MatchStructuredDimOp::matchOperation(
     transform::TransformState &state) {
   auto linalgOp = cast<linalg::LinalgOp>(current);
   SmallVector<int64_t> dimensions;
   DiagnosedSilenceableFailure diag = getDimensionsFor(linalgOp, dimensions);
   if (!diag.succeeded())
     return diag;
 
   // If asked to check for the kind of dimension, perform the check.
-  if (getParallel() || getReduction()) {
+  if (isParallel() || isReduction()) {
     SmallVector<unsigned> reference;
-    if (getParallel())
+    if (isParallel())
       linalgOp.getParallelDims(reference);
-    else if (getReduction())
+    else if (isReduction())
       linalgOp.getReductionDims(reference);
 
     DiagnosedSilenceableFailure diag =
         containsAll(reference, dimensions, getLoc(),
-                    getParallel() ? "expects dimension #{0} to be parallel"
-                                  : "expects dimension #{0} to be reduction");
+                    isParallel() ? "expects dimension #{0} to be parallel"
     if (!diag.succeeded())
       return diag;
   }
 
   // If not capturing, we are done here.
   if (!getResult())
     return diag;
 
   SmallVector<int64_t, 4> ranges = linalgOp.getStaticLoopRanges();
   Builder builder(current);
   SmallVector<Attribute> captured = llvm::to_vector(
       llvm::map_range(dimensions, [&](int64_t dim) -> Attribute {
         return builder.getI64IntegerAttr(ranges[dim]);
       }));
   results.setParams(cast<OpResult>(getResult()), captured);
   return DiagnosedSilenceableFailure::success();
 }
 
 DiagnosedSilenceableFailure transform::MatchStructuredDimOp::getDimensionsFor(
     linalg::LinalgOp op, SmallVectorImpl<int64_t> &dims) {
-  DiagnosedSilenceableFailure diag =
-      expandTargetSpecification(getLoc(), getIsAll(), getIsInverted(),
-                                getRawDimList(), op.getNumLoops(), dims);
+  DiagnosedSilenceableFailure diag = expandTargetSpecification(
+      getLoc(), isAll(), isInverted(), getRawDimList(), op.getNumLoops(), dims);
   if (diag.isSilenceableFailure()) {
     diag.attachNote(op->getLoc())
         << "while considering dimensions of this payload operation";
   }
   return diag;
 }
 
 LogicalResult transform::MatchStructuredDimOp::verify() {
-  if (getParallel() && getReduction()) {
+  if (isParallel() && isReduction()) {
     return emitOpError() << "cannot request the same dimension to be both "
                             "parallel and reduction";
   }
   return verifyStructuredTransformDimsOp(getOperation(), getRawDimList(),
-                                         getIsInverted(), getIsAll());
+                                         isInverted(), isAll());
 }
 
 //===----------------------------------------------------------------------===//
 // MatchStructuredElementalBitwidthOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure
 transform::MatchStructuredElementalBitwidthOp::matchValue(
DiagnosedSilenceableFailure transform::MatchStructuredInputOp::matchOperation(
   if (!diag.succeeded())
     return diag;
 
   SmallVector<MappedValue> operandMapping;
   operandMapping.reserve(positions.size());
   for (int64_t position : positions) {
     AffineMap indexingMap =
         linalgOp.getMatchingIndexingMap(linalgOp.getDpsInputOperand(position));
-    if (getPermutation() && !indexingMap.isPermutation()) {
+    if (isPermutation() && !indexingMap.isPermutation()) {
       return emitSilenceableError() << "the indexing map for input #"
                                     << position << " is not a permutation";
     }
-    if (getProjectedPermutation() && !indexingMap.isProjectedPermutation()) {
+    if (isProjectedPermutation() && !indexingMap.isProjectedPermutation()) {
       return emitSilenceableError()
              << "the indexing map for input #" << position
              << " is not a projected permutation";
     }
 
     // If capture not requested, skip it.
     if (!getResult())
       continue;
DiagnosedSilenceableFailure transform::MatchStructuredInputOp::matchOperation(
   if (getResult())
     results.setMappedValues(cast<OpResult>(getResult()), operandMapping);
   return DiagnosedSilenceableFailure::success();
 }
 
 DiagnosedSilenceableFailure transform::MatchStructuredInputOp::getPositionsFor(
     linalg::LinalgOp op, SmallVectorImpl<int64_t> &positions) {
   DiagnosedSilenceableFailure diag = expandTargetSpecification(
-      getLoc(), getIsAll(), getIsInverted(), getRawPositionList(),
+      getLoc(), isAll(), isInverted(), getRawPositionList(),
       op.getNumDpsInputs(), positions);
   if (diag.isSilenceableFailure()) {
     diag.attachNote(op->getLoc())
         << "while considering DPS inputs of this payload operation";
   }
   return diag;
 }
 
 /// Verifies a matcher op for structured input or output, specifically the
 /// attributes specifying the operand positions.
 template <typename OpTy>
 LogicalResult verifyStructuredOperandOp(OpTy op) {
-  if (op.getPermutation() && op.getProjectedPermutation()) {
+  if (op.isPermutation() && op.isProjectedPermutation()) {
     return op.emitOpError()
            << op.getPermutationAttrName() << " and "
            << op.getProjectedPermutationAttrName() << " are mutually exclusive";
   }
   if (op.getRawPositionList().size() > 1 && op.getResult()) {
     return op.emitOpError()
            << "cannot bind multiple inputs/inits to the same value";
   }
 
   return success();
 }
 
 LogicalResult transform::MatchStructuredInputOp::verify() {
   if (failed(verifyStructuredOperandOp(*this)))
     return failure();
   return verifyStructuredTransformDimsOp(getOperation(), getRawPositionList(),
-                                         getIsInverted(), getIsAll());
+                                         isInverted(), isAll());
 }
 
 //===----------------------------------------------------------------------===//
 // MatchStructuredInitOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure transform::MatchStructuredInitOp::matchOperation(
     Operation *current, transform::TransformResults &results,
     transform::TransformState &state) {
   auto linalgOp = cast<linalg::LinalgOp>(current);
   SmallVector<int64_t> positions;
   DiagnosedSilenceableFailure diag = getPositionsFor(linalgOp, positions);
   if (!diag.succeeded())
     return diag;
 
   SmallVector<MappedValue> operandMapping;
   operandMapping.reserve(positions.size());
   for (int64_t position : positions) {
     AffineMap indexingMap =
         linalgOp.getMatchingIndexingMap(linalgOp.getDpsInitOperand(position));
-    if (getPermutation() && !indexingMap.isPermutation()) {
+    if (isPermutation() && !indexingMap.isPermutation()) {
       return emitSilenceableError() << "the indexing map for output(init) #"
                                     << position << " is not a permutation";
     }
-    if (getProjectedPermutation() && !indexingMap.isProjectedPermutation()) {
+    if (isProjectedPermutation() && !indexingMap.isProjectedPermutation()) {
       return emitSilenceableError() << "the indexing map for output(init) #"
                                     << position << " is not a permutation";
     }
 
     // If capture not requested, skip it.
     if (!getResult())
       continue;
 
DiagnosedSilenceableFailure transform::MatchStructuredInitOp::matchOperation(
   if (getResult())
     results.setMappedValues(cast<OpResult>(getResult()), operandMapping);
   return DiagnosedSilenceableFailure::success();
 }
 
 DiagnosedSilenceableFailure transform::MatchStructuredInitOp::getPositionsFor(
     linalg::LinalgOp op, SmallVectorImpl<int64_t> &positions) {
   DiagnosedSilenceableFailure diag = expandTargetSpecification(
-      getLoc(), getIsAll(), getIsInverted(), getRawPositionList(),
+      getLoc(), isAll(), isInverted(), getRawPositionList(),
       op.getNumDpsInits(), positions);
   if (diag.isSilenceableFailure()) {
     diag.attachNote(op->getLoc())
         << "while considering DPS inits (outputs) of this payload operation";
   }
   return diag;
 }
 
 LogicalResult transform::MatchStructuredInitOp::verify() {
   if (failed(verifyStructuredOperandOp(*this)))
     return failure();
   return verifyStructuredTransformDimsOp(getOperation(), getRawPositionList(),
-                                         getIsInverted(), getIsAll());
+                                         isInverted(), isAll());
 }
 
 //===----------------------------------------------------------------------===//
 // MatchStructuredNumInputsOp
 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure
 transform::MatchStructuredNumInputsOp::matchOperation(
if (isa<TransformValueHandleTypeInterface>(getResult().getType())) {
     return DiagnosedSilenceableFailure::success();
   }
 
   if (result.getUsers().empty()) {
     return emitSilenceableError()
            << "no users of the result #" << getPosition();
   }
   Operation *firstUser = *result.getUsers().begin();
-  if (getAny()) {
+  if (isAny()) {
     results.set(cast<OpResult>(getResult()), {firstUser});
     return DiagnosedSilenceableFailure::success();
   }
-  if (getSingle()) {
+  if (isSingle()) {
     if (!llvm::hasSingleElement(result.getUsers())) {
       return emitSilenceableError()
              << "more than one result user with single user requested";
     }
     results.set(cast<OpResult>(getResult()), {firstUser});
     return DiagnosedSilenceableFailure::success();
   }
 
if (position >= op.getNumDpsInits() || position < 0) {
     return emitSilenceableError()
            << "position " << rawPosition
            << " overflows the number of results(ints) of the payload operation";
   }
   return DiagnosedSilenceableFailure::success();
 }
 
 LogicalResult transform::MatchStructuredResultOp::verify() {
-  if ((getAny() || getSingle()) ^
+  if ((isAny() || isSingle()) ^
       isa<TransformHandleTypeInterface>(getResult().getType())) {
     return emitOpError() << "expects either the any/single keyword or the type "
                             "value handle result type";
   }
-  if (getAny() && getSingle()) {
+  if (isAny() && isSingle()) {
     return emitOpError() << "'any' and 'single' are mutually exclusive";
   }
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // MatchStructuredYieldOp
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/Linalg/TransformOps/LinalgTransformOps.cpp

 DiagnosedSilenceableFailure
 transform::PromoteOp::applyToOne(LinalgOp target,
                                  transform::ApplyToEachResultList &results,
                                  transform::TransformState &state) {
   LinalgPromotionOptions promotionOptions;
   if (!getOperandsToPromote().empty())
     promotionOptions = promotionOptions.setOperandsToPromote(
         extractFromI64ArrayAttr(getOperandsToPromote()));
-  if (getUseFullTilesByDefault())
+  if (isUseFullTilesByDefault())
     promotionOptions = promotionOptions.setUseFullTileBuffersByDefault(
-        getUseFullTilesByDefault());
-  if (getUseAlloca())
-    promotionOptions = promotionOptions.setUseAlloca(getUseAlloca());
+        isUseFullTilesByDefault());
+  if (isUseAlloca())
+    promotionOptions = promotionOptions.setUseAlloca(isUseAlloca());
   if (!getUseFullTileBuffers().empty())
     promotionOptions = promotionOptions.setUseFullTileBuffers(
         llvm::to_vector(getUseFullTileBuffers().getAsValueRange<BoolAttr>()));
   if (getAlignment().has_value())
     promotionOptions = promotionOptions.setAlignment(*getAlignment());
 
   if (getMapping().has_value()) {
     // The mapping should only contain an element

 }
 
 DiagnosedSilenceableFailure transform::SplitReductionOp::applyToOne(
     LinalgOp target, transform::ApplyToEachResultList &results,
     transform::TransformState &state) {
   ControlSplitReductionFn splitFn = [&](LinalgOp) {
     return linalg::SplitReductionOptions{int64_t(getSplitFactor()),
                                          unsigned(getInsertSplitDimension()),
-                                         bool(getInnerParallel())};
+                                         bool(isInnerParallel())};
   };
   TrackingListener listener(state, *this);
   IRRewriter rewriter(getContext(), &listener);
   rewriter.setInsertionPoint(target);
   FailureOr<SplitReductionResult> splitResult =
-      (getUseScalingAlgorithm())
-          ? splitReductionByScaling(rewriter, target, splitFn, getUseAlloc())
-          : splitReduction(rewriter, target, splitFn, getUseAlloc());
+      (isUseScalingAlgorithm())
+          ? splitReductionByScaling(rewriter, target, splitFn, isUseAlloc())
+          : splitReduction(rewriter, target, splitFn, isUseAlloc());
   if (failed(splitResult))
     return emitDefaultDefiniteFailure(target);
 
   results.push_back(splitResult->initOrAlloc);
   results.push_back(splitResult->fillOp);
   results.push_back(splitResult->splitLinalgOp);
   results.push_back(splitResult->resultCombiningLinalgOp);
   return DiagnosedSilenceableFailure::success();
transform::VectorizeOp::applyToOne(Operation *target,
   if (!target->hasTrait<OpTrait::IsIsolatedFromAbove>()) {
     auto diag = this->emitOpError("requires isolated-from-above targets");
     diag.attachNote(target->getLoc()) << "non-isolated target";
     return DiagnosedSilenceableFailure::definiteFailure();
   }
 
   MLIRContext *ctx = getContext();
   RewritePatternSet patterns(ctx);
-  patterns.add<VectorizationPattern>(ctx, getVectorizeNdExtract());
+  patterns.add<VectorizationPattern>(ctx, isVectorizeNdExtract());
 
-  if (!getDisableTransferPermutationMapLoweringPatterns())
+  if (!isDisableTransferPermutationMapLoweringPatterns())
     vector::populateVectorTransferPermutationMapLoweringPatterns(patterns);
 
-  if (!getDisableMultiReductionToContractPatterns())
+  if (!isDisableMultiReductionToContractPatterns())
     vector::populateVectorReductionToContractPatterns(patterns);
 
   patterns.add<linalg::LinalgCopyVTRForwardingPattern,
                linalg::LinalgCopyVTWForwardingPattern>(ctx,
                                                        /*benefit=*/2);
   vector::TransferReadOp::getCanonicalizationPatterns(patterns, ctx);
   vector::TransferWriteOp::getCanonicalizationPatterns(patterns, ctx);
   vector::populateVectorTransferTensorSliceTransforms(patterns);
 
   patterns.add<CopyVectorizationPattern>(ctx);
 
-  if (getVectorizePadding())
+  if (isVectorizePadding())
     linalg::populatePadOpVectorizationPatterns(patterns);
 
   TrackingListener listener(state, *this);
   GreedyRewriteConfig config;
   config.listener = &listener;
   if (failed(applyPatternsAndFoldGreedily(target, std::move(patterns), config)))
     return emitDefaultDefiniteFailure(target);
 
for (Operation *target : targets) {
 
     auto linalgOp = dyn_cast<LinalgOp>(target);
     if (!linalgOp) {
       return mlir::emitSilenceableFailure(target->getLoc())
              << "cannot vectorize non-Linalg op";
     }
 
     if (failed(linalg::vectorize(rewriter, linalgOp, vectorSizes,
-                                 getVectorizeNdExtract()))) {
+                                 isVectorizeNdExtract()))) {
       return mlir::emitSilenceableFailure(target->getLoc())
              << "failed to vectorize linalg op";
     }
   }
 
   return DiagnosedSilenceableFailure::success();
 }
 

mlir/lib/Dialect/Linalg/Transforms/ConvertToDestinationStyle.cpp

return rewriter.notifyMatchFailure(
         padOp, "failed to reify tensor.pad op result shape");
   SmallVector<Value> dynamicSizes;
   for (int64_t i = 0; i < resultType.getRank(); ++i)
     if (resultType.isDynamicDim(i))
       dynamicSizes.push_back(reifiedShape[0][i].get<Value>());
 
   // If the `padOp` has a nofold attribute and all paddings are known to be 0,
   // explicitly insert a `linalg.copy`.
-  if (padOp.getNofoldAttr() &&
+  if (padOp.isNofoldAttr() &&
       llvm::all_of(padOp.getMixedLowPad(), isZeroIndex) &&
       llvm::all_of(padOp.getMixedHighPad(), isZeroIndex)) {
     using bufferization::AllocTensorOp;
     Value allocated =
         rewriter.create<AllocTensorOp>(loc, resultType, dynamicSizes);
     auto copyOp = rewriter.replaceOpWithNewOp<linalg::CopyOp>(
         padOp, padOp.getSource(), allocated);
     return copyOp.getOperation();

mlir/lib/Dialect/Linalg/Transforms/DataLayoutPropagation.cpp

LogicalResult matchAndRewrite(tensor::PadOp padOp,
     }
     // Add zero padding for the point loops.
     size_t pointLoopsSize = innerDimsPos.size();
     lowPad.append(pointLoopsSize, rewriter.getIndexAttr(0));
     highPad.append(pointLoopsSize, rewriter.getIndexAttr(0));
 
     auto newPadOp = rewriter.create<tensor::PadOp>(
         loc, /*result=*/Type(), unpackOp.getSource(), lowPad, highPad,
-        paddingVal, padOp.getNofold());
+        paddingVal, padOp.isNofold());
 
     // Inject the tensor.unpack right after the packed padOp.
     Value outputUnPack = rewriter.create<tensor::EmptyOp>(
         loc, padOp.getResultType().getShape(),
         padOp.getResultType().getElementType());
 
     Value replacement = rewriter.create<tensor::UnPackOp>(
         loc, newPadOp.getResult(), outputUnPack, innerDimsPos,

mlir/lib/Dialect/MLProgram/IR/MLProgramOps.cpp

function_interface_impl::printFunctionOp(
       getArgAttrsAttrName(), getResAttrsAttrName());
 }
 
 //===----------------------------------------------------------------------===//
 // GlobalOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult GlobalOp::verify() {
-  if (!getIsMutable() && !getValue())
+  if (!isGlobalMutable() && !getValue())
     return emitOpError() << "immutable global must have an initial value";
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // GlobalLoadOp
 //===----------------------------------------------------------------------===//
 

 }
 
 LogicalResult
 GlobalLoadConstOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
   GlobalOp referrent = getGlobalOp(symbolTable);
   if (!referrent)
     return emitOpError() << "undefined global: " << getGlobal();
 
-  if (referrent.getIsMutable())
+  if (referrent.isGlobalMutable())
     return emitOpError() << "cannot load as const from mutable global "
                          << getGlobal();
 
   if (referrent.getType() != getResult().getType())
     return emitOpError() << "cannot load from global typed "
                          << referrent.getType() << " as "
                          << getResult().getType();
 

 }
 
 LogicalResult
 GlobalStoreOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
   GlobalOp referrent = getGlobalOp(symbolTable);
   if (!referrent)
     return emitOpError() << "undefined global: " << getGlobal();
 
-  if (!referrent.getIsMutable()) {
+  if (!referrent.isGlobalMutable()) {
     return emitOpError() << "cannot store to an immutable global "
                          << getGlobal();
   }
 
   if (referrent.getType() != getValue().getType()) {
     return emitOpError() << "cannot store to a global typed "
                          << referrent.getType() << " from "
                          << getValue().getType();

 }
 
 LogicalResult
 GlobalStoreGraphOp::verifySymbolUses(SymbolTableCollection &symbolTable) {
   GlobalOp referrent = getGlobalOp(symbolTable);
   if (!referrent)
     return emitOpError() << "undefined global: " << getGlobal();
 
-  if (!referrent.getIsMutable()) {
+  if (!referrent.isGlobalMutable()) {
     return emitOpError() << "cannot store to an immutable global "
                          << getGlobal();
   }
 
   if (referrent.getType() != getValue().getType()) {
     return emitOpError() << "cannot store to a global typed "
                          << referrent.getType() << " from "
                          << getValue().getType();

mlir/lib/Dialect/MemRef/IR/MemRefOps.cpp

LogicalResult GlobalOp::verify() {
   }
 
   // TODO: verify visibility for declarations.
   return success();
 }
 
 ElementsAttr GlobalOp::getConstantInitValue() {
   auto initVal = getInitialValue();
-  if (getConstant() && initVal.has_value())
+  if (isConstant() && initVal.has_value())
     return llvm::cast<ElementsAttr>(initVal.value());
   return {};
 }
 
 //===----------------------------------------------------------------------===//
 // GetGlobalOp
 //===----------------------------------------------------------------------===//
 

mlir/lib/Dialect/MemRef/TransformOps/MemRefTransformOps.cpp

for (Operation *user : target->getUsers()) {
       }
     }
     if (!canApplyMultiBuffer) {
       LLVM_DEBUG(DBGS() << "--cannot apply multibuffering -> Skip\n";);
       continue;
     }
 
     auto newBuffer =
-        memref::multiBuffer(rewriter, target, getFactor(), getSkipAnalysis());
+        memref::multiBuffer(rewriter, target, getFactor(), isSkipAnalysis());
 
     if (failed(newBuffer)) {
       LLVM_DEBUG(DBGS() << "--op failed to multibuffer\n";);
       return emitSilenceableFailure(target->getLoc())
              << "op failed to multibuffer";
     }
 
     results.push_back(*newBuffer);

mlir/lib/Dialect/MemRef/Transforms/FoldMemRefAliasOps.cpp

llvm::TypeSwitch<Operation *, void>(loadOp)
             AffineMapAttr::get(expandDimsToRank(
                 op.getPermutationMap(), subViewOp.getSourceType().getRank(),
                 subViewOp.getDroppedDims())),
             op.getPadding(), /*mask=*/Value(), op.getInBoundsAttr());
       })
       .Case([&](gpu::SubgroupMmaLoadMatrixOp op) {
         rewriter.replaceOpWithNewOp<gpu::SubgroupMmaLoadMatrixOp>(
             op, op.getType(), subViewOp.getSource(), sourceIndices,
-            op.getLeadDimension(), op.getTransposeAttr());
+            op.getLeadDimension(), op.isTransposeAttr());
       })
       .Default([](Operation *) { llvm_unreachable("unexpected operation."); });
   return success();
 }
 
 template <typename OpTy>
 LogicalResult LoadOpOfExpandShapeOpFolder<OpTy>::matchAndRewrite(
     OpTy loadOp, PatternRewriter &rewriter) const {
llvm::TypeSwitch<Operation *, void>(storeOp)
             AffineMapAttr::get(expandDimsToRank(
                 op.getPermutationMap(), subViewOp.getSourceType().getRank(),
                 subViewOp.getDroppedDims())),
             op.getInBoundsAttr());
       })
       .Case([&](gpu::SubgroupMmaStoreMatrixOp op) {
         rewriter.replaceOpWithNewOp<gpu::SubgroupMmaStoreMatrixOp>(
             op, op.getSrc(), subViewOp.getSource(), sourceIndices,
-            op.getLeadDimension(), op.getTransposeAttr());
+            op.getLeadDimension(), op.isTransposeAttr());
       })
       .Default([](Operation *) { llvm_unreachable("unexpected operation."); });
   return success();
 }
 
 template <typename OpTy>
 LogicalResult StoreOpOfExpandShapeOpFolder<OpTy>::matchAndRewrite(
     OpTy storeOp, PatternRewriter &rewriter) const {
LogicalResult NvgpuAsyncCopyOpSubViewOpFolder::matchAndRewrite(
   // Replace the copy op with a new copy op that uses the source and destination
   // of the subview.
   rewriter.replaceOpWithNewOp<nvgpu::DeviceAsyncCopyOp>(
       copyOp, nvgpu::DeviceAsyncTokenType::get(copyOp.getContext()),
       (dstSubViewOp ? dstSubViewOp.getSource() : copyOp.getDst()),
       foldedDstIndices,
       (srcSubViewOp ? srcSubViewOp.getSource() : copyOp.getSrc()),
       foldedSrcIndices, copyOp.getDstElements(), copyOp.getSrcElements(),
-      copyOp.getBypassL1Attr());
+      copyOp.isBypassL1Attr());
 
   return success();
 }
 
 void memref::populateFoldMemRefAliasOpPatterns(RewritePatternSet &patterns) {
   patterns.add<LoadOpOfSubViewOpFolder<affine::AffineLoadOp>,
                LoadOpOfSubViewOpFolder<memref::LoadOp>,
                LoadOpOfSubViewOpFolder<vector::TransferReadOp>,

mlir/lib/Dialect/OpenACC/IR/OpenACC.cpp

 void printVectorClause(OpAsmPrinter &p, Operation *op, Value vectorLength,
                        Type vectorLengthType, UnitAttr hasVector) {
   if (vectorLength)
     p << "(" << vectorLength << " : " << vectorLengthType << ")";
 }
 
 LogicalResult acc::LoopOp::verify() {
   // auto, independent and seq attribute are mutually exclusive.
-  if ((getAuto_() && (getIndependent() || getSeq())) ||
-      (getIndependent() && getSeq())) {
+  if ((isAuto_() && (isIndependent() || isSeq())) ||
+      (isIndependent() && isSeq())) {
     return emitError() << "only one of \"" << acc::LoopOp::getAutoAttrStrName()
                        << "\", " << getIndependentAttrName() << ", "
-                       << getSeqAttrName()
-                       << " can be present at the same time";
+                       << getSeqAttrName() << " can be present at the same time";
   }
 
   // Gang, worker and vector are incompatible with seq.
-  if (getSeq() && (getHasGang() || getHasWorker() || getHasVector()))
+  if (isSeq() && (isHasGang() || isHasWorker() || isHasVector()))
     return emitError("gang, worker or vector cannot appear with the seq attr");
 
   // Check non-empty body().
   if (getRegion().empty())
     return emitError("expected non-empty body.");
 
   return success();
 }
LogicalResult acc::ExitDataOp::verify() {
   // At least one copyout, delete, or detach clause must appear on an exit data
   // directive.
   if (getDataClauseOperands().empty())
     return emitError("at least one operand must be present in dataOperands on "
                      "the exit data operation");
 
   // The async attribute represent the async clause without value. Therefore the
   // attribute and operand cannot appear at the same time.
-  if (getAsyncOperand() && getAsync())
+  if (getAsyncOperand() && isAsync())
     return emitError("async attribute cannot appear with asyncOperand");
 
   // The wait attribute represent the wait clause without values. Therefore the
   // attribute and operands cannot appear at the same time.
-  if (!getWaitOperands().empty() && getWait())
+  if (!getWaitOperands().empty() && isWait())
     return emitError("wait attribute cannot appear with waitOperands");
 
   if (getWaitDevnum() && getWaitOperands().empty())
     return emitError("wait_devnum cannot appear without waitOperands");
 
   return success();
 }
 
LogicalResult acc::EnterDataOp::verify() {
   // At least one copyin, create, or attach clause must appear on an enter data
   // directive.
   if (getDataClauseOperands().empty())
     return emitError("at least one operand must be present in dataOperands on "
                      "the enter data operation");
 
   // The async attribute represent the async clause without value. Therefore the
   // attribute and operand cannot appear at the same time.
-  if (getAsyncOperand() && getAsync())
+  if (getAsyncOperand() && isAsync())
     return emitError("async attribute cannot appear with asyncOperand");
 
   // The wait attribute represent the wait clause without values. Therefore the
   // attribute and operands cannot appear at the same time.
-  if (!getWaitOperands().empty() && getWait())
+  if (!getWaitOperands().empty() && isWait())
     return emitError("wait attribute cannot appear with waitOperands");
 
   if (getWaitDevnum() && getWaitOperands().empty())
     return emitError("wait_devnum cannot appear without waitOperands");
 
   for (mlir::Value operand : getDataClauseOperands())
     if (!mlir::isa<acc::AttachOp, acc::CreateOp, acc::CopyinOp>(
             operand.getDefiningOp()))

 
 LogicalResult acc::UpdateOp::verify() {
   // At least one of host or device should have a value.
   if (getDataClauseOperands().empty())
     return emitError("at least one value must be present in dataOperands");
 
   // The async attribute represent the async clause without value. Therefore the
   // attribute and operand cannot appear at the same time.
-  if (getAsyncOperand() && getAsync())
+  if (getAsyncOperand() && isAsync())
     return emitError("async attribute cannot appear with asyncOperand");
 
   // The wait attribute represent the wait clause without values. Therefore the
   // attribute and operands cannot appear at the same time.
-  if (!getWaitOperands().empty() && getWait())
+  if (!getWaitOperands().empty() && isWait())
     return emitError("wait attribute cannot appear with waitOperands");
 
   if (getWaitDevnum() && getWaitOperands().empty())
     return emitError("wait_devnum cannot appear without waitOperands");
 
   for (mlir::Value operand : getDataClauseOperands())
     if (!mlir::isa<acc::UpdateDeviceOp, acc::UpdateHostOp, acc::GetDevicePtrOp>(
             operand.getDefiningOp()))

 
 //===----------------------------------------------------------------------===//
 // WaitOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult acc::WaitOp::verify() {
   // The async attribute represent the async clause without value. Therefore the
   // attribute and operand cannot appear at the same time.
-  if (getAsyncOperand() && getAsync())
+  if (getAsyncOperand() && isAsync())
     return emitError("async attribute cannot appear with asyncOperand");
 
   if (getWaitDevnum() && getWaitOperands().empty())
     return emitError("wait_devnum cannot appear without waitOperands");
 
   return success();
 }
 
 #define GET_OP_CLASSES
 #include "mlir/Dialect/OpenACC/OpenACCOps.cpp.inc"
 
 #define GET_ATTRDEF_CLASSES
 #include "mlir/Dialect/OpenACC/OpenACCOpsAttributes.cpp.inc"
 
 #define GET_TYPEDEF_CLASSES
 #include "mlir/Dialect/OpenACC/OpenACCOpsTypes.cpp.inc"

mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp

if (getAllocateVars().size() != getAllocatorsVars().size())
     return emitError(
         "expected equal sizes for allocate and allocator variables");
   if (failed(
           verifyReductionVarList(*this, getReductions(), getReductionVars())) ||
       failed(verifyReductionVarList(*this, getInReductions(),
                                     getInReductionVars())))
     return failure();
 
-  if (!getReductionVars().empty() && getNogroup())
+  if (!getReductionVars().empty() && isNogroup())
     return emitError("if a reduction clause is present on the taskloop "
                      "directive, the nogroup clause must not be specified");
   for (auto var : getReductionVars()) {
     if (llvm::is_contained(getInReductionVars(), var))
       return emitError("the same list item cannot appear in both a reduction "
                        "and an in_reduction clause");
   }
 
return emitOpError() << "number of variables in depend clause does not "
                          << "match number of iteration variables in the "
                          << "doacross loop";
 
   return success();
 }
 
 LogicalResult OrderedRegionOp::verify() {
   // TODO: The code generation for ordered simd directive is not supported yet.
-  if (getSimd())
+  if (isSimd())
     return failure();
 
   if (auto container = (*this)->getParentOfType<WsLoopOp>()) {
     if (!container.getOrderedValAttr() ||
         container.getOrderedValAttr().getInt() != 0)
       return emitOpError() << "ordered region must be closely nested inside "
                            << "a worksharing-loop region with an ordered "
                            << "clause without parameter present";
if ((cct == ClauseCancellationConstructType::Parallel) &&
     return emitOpError() << "cancel parallel must appear "
                          << "inside a parallel region";
   }
   if (cct == ClauseCancellationConstructType::Loop) {
     if (!isa<WsLoopOp>(parentOp)) {
       return emitOpError() << "cancel loop must appear "
                            << "inside a worksharing-loop region";
     }
-    if (cast<WsLoopOp>(parentOp).getNowaitAttr()) {
+    if (cast<WsLoopOp>(parentOp).isNowaitAttr()) {
       return emitError() << "A worksharing construct that is canceled "
                          << "must not have a nowait clause";
     }
     if (cast<WsLoopOp>(parentOp).getOrderedValAttr()) {
       return emitError() << "A worksharing construct that is canceled "
                          << "must not have an ordered clause";
     }
 
   } else if (cct == ClauseCancellationConstructType::Sections) {
     if (!(isa<SectionsOp>(parentOp) || isa<SectionOp>(parentOp))) {
       return emitOpError() << "cancel sections must appear "
                            << "inside a sections region";
     }
     if (isa_and_nonnull<SectionsOp>(parentOp->getParentOp()) &&
-        cast<SectionsOp>(parentOp->getParentOp()).getNowaitAttr()) {
+        cast<SectionsOp>(parentOp->getParentOp()).isNowaitAttr()) {
       return emitError() << "A sections construct that is canceled "
                          << "must not have a nowait clause";
     }
   }
   // TODO : Add more when we support taskgroup.
   return success();
 }
 //===----------------------------------------------------------------------===//

mlir/lib/Dialect/PDLInterp/IR/PDLInterp.cpp

static LogicalResult verifySwitchOp(OpT op) {
   return success();
 }
 
 //===----------------------------------------------------------------------===//
 // pdl_interp::CreateOperationOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult CreateOperationOp::verify() {
-  if (!getInferredResultTypes())
+  if (!isInferredResultTypes())
     return success();
   if (!getInputResultTypes().empty()) {
     return emitOpError("with inferred results cannot also have "
                        "explicit result types");
   }
   OperationName opName(getName(), getContext());
   if (!opName.hasInterface<InferTypeOpInterface>()) {
     return emitOpError()

mlir/lib/Dialect/SCF/TransformOps/SCFTransformOps.cpp

 DiagnosedSilenceableFailure transform::TakeAssumedBranchOp::applyToOne(
     scf::IfOp ifOp, transform::ApplyToEachResultList &results,
     transform::TransformState &state) {
   TrackingListener listener(state, *this);
   IRRewriter rewriter(ifOp->getContext(), &listener);
   rewriter.setInsertionPoint(ifOp);
 
   Region &region =
-      getTakeElseBranch() ? ifOp.getElseRegion() : ifOp.getThenRegion();
+      isTakeElseBranch() ? ifOp.getElseRegion() : ifOp.getThenRegion();
   if (!llvm::hasSingleElement(region)) {
     return emitDefiniteFailure()
            << "requires an scf.if op with a single-block "
-           << ((getTakeElseBranch()) ? "`else`" : "`then`") << " region";
+           << ((isTakeElseBranch()) ? "`else`" : "`then`") << " region";
   }
   replaceOpWithRegion(rewriter, ifOp, region);
   return DiagnosedSilenceableFailure::success();
 }
 
 void transform::TakeAssumedBranchOp::getEffects(
     SmallVectorImpl<MemoryEffects::EffectInstance> &effects) {
   onlyReadsHandle(getTarget(), effects);

mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp

LogicalResult BinaryOp::verify() {
   // non-empty region.
   if (!overlap.empty()) {
     RETURN_FAILURE_IF_FAILED(verifyNumBlockArgs(
         this, overlap, "overlap", TypeRange{leftType, rightType}, outputType))
   }
   if (!left.empty()) {
     RETURN_FAILURE_IF_FAILED(
         verifyNumBlockArgs(this, left, "left", TypeRange{leftType}, outputType))
-  } else if (getLeftIdentity()) {
+  } else if (isLeftIdentity()) {
     if (leftType != outputType)
       return emitError("left=identity requires first argument to have the same "
                        "type as the output");
   }
   if (!right.empty()) {
     RETURN_FAILURE_IF_FAILED(verifyNumBlockArgs(
         this, right, "right", TypeRange{rightType}, outputType))
-  } else if (getRightIdentity()) {
+  } else if (isRightIdentity()) {
     if (rightType != outputType)
       return emitError("right=identity requires second argument to have the "
                        "same type as the output");
   }
   return success();
 }
 
 LogicalResult UnaryOp::verify() {

mlir/lib/Dialect/SparseTensor/Transforms/SparseBufferRewriting.cpp

LogicalResult matchAndRewrite(PushBackOp op,
     Value size = op.getCurSize();
     Value value = op.getValue();
 
     Value n = op.getN() ? op.getN() : constantIndex(rewriter, loc, 1);
     Value newSize = rewriter.create<arith::AddIOp>(loc, size, n);
     auto nValue = dyn_cast_or_null<arith::ConstantIndexOp>(n.getDefiningOp());
     bool nIsOne = (nValue && nValue.value() == 1);
 
-    if (!op.getInbounds()) {
+    if (!op.isInbounds()) {
       Value cond = rewriter.create<arith::CmpIOp>(
           loc, arith::CmpIPredicate::ugt, newSize, capacity);
 
       Value c2 = constantIndex(rewriter, loc, 2);
       auto bufferType =
           MemRefType::get({ShapedType::kDynamic}, value.getType());
       scf::IfOp ifOp = rewriter.create<scf::IfOp>(loc, bufferType, cond,
                                                   /*else=*/true);

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorCodegen.cpp

 public:
   using OpConversionPattern::OpConversionPattern;
   LogicalResult
   matchAndRewrite(LoadOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     // Prepare descriptor.
     auto desc = getDescriptorFromTensorTuple(adaptor.getTensor());
     // Generate optional insertion finalization code.
-    if (op.getHasInserts())
+    if (op.isHasInserts())
       genEndInsert(rewriter, op.getLoc(), desc);
     // Replace operation with resulting memrefs.
     rewriter.replaceOp(op, genTuple(rewriter, op.getLoc(), desc));
     return success();
   }
 };
 
 /// Sparse codegen rule for the expand op.

mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp

 
 /// Sparse conversion rule for tensor rematerialization.
 class SparseTensorLoadConverter : public OpConversionPattern<LoadOp> {
 public:
   using OpConversionPattern::OpConversionPattern;
   LogicalResult
   matchAndRewrite(LoadOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
-    if (op.getHasInserts()) {
+    if (op.isHasInserts()) {
       // Finalize any pending insertions.
       StringRef name = "endInsert";
       createFuncCall(rewriter, op->getLoc(), name, {}, adaptor.getOperands(),
                      EmitCInterface::Off);
     }
     rewriter.replaceOp(op, adaptor.getOperands());
     return success();
   }

mlir/lib/Dialect/SparseTensor/Utils/Merger.cpp

//     x  | left(x) | overlap(x,y) |
         leftYield = leftBlock.getTerminator();
       }
       // Right Region.
       Operation *rightYield = nullptr;
       if (!rightRegion.empty()) {
         Block &rightBlock = rightRegion.front();
         rightYield = rightBlock.getTerminator();
       }
-      bool includeLeft = binop.getLeftIdentity() || !leftRegion.empty();
-      bool includeRight = binop.getRightIdentity() || !rightRegion.empty();
+      bool includeLeft = binop.isLeftIdentity() || !leftRegion.empty();
+      bool includeRight = binop.isRightIdentity() || !rightRegion.empty();
       return combiSet(TensorExp::Kind::kBinary, child0, child1, binop,
                       includeLeft, TensorExp::Kind::kBinaryBranch, leftYield,
                       includeRight, TensorExp::Kind::kBinaryBranch, rightYield);
     }
   case TensorExp::Kind::kReduce:
     // A custom reduce operation.
     {
       const ExprId e0 = expr.children.e0;
if (x.has_value() && y.has_value()) {
       if (isa<arith::ShRSIOp>(def) && isInvariant(e1))
         return addExp(TensorExp::Kind::kShrS, e0, e1);
       if (isa<arith::ShRUIOp>(def) && isInvariant(e1))
         return addExp(TensorExp::Kind::kShrU, e0, e1);
       if (isa<arith::ShLIOp>(def) && isInvariant(e1))
         return addExp(TensorExp::Kind::kShlI, e0, e1);
       if (auto binop = dyn_cast<sparse_tensor::BinaryOp>(def)) {
         if (isAdmissibleBranch(binop, binop.getOverlapRegion()) &&
-            (binop.getLeftIdentity() ||
+            (binop.isLeftIdentity() ||
              isAdmissibleBranch(binop, binop.getLeftRegion())) &&
-            (binop.getRightIdentity() ||
+            (binop.isRightIdentity() ||
              isAdmissibleBranch(binop, binop.getRightRegion())))
           return addExp(TensorExp::Kind::kBinary, e0, e1, def);
       }
     }
   }
   // Construct ternary operations if subexpressions can be built.
   if (def->getNumOperands() == 3) {
     const auto x = buildTensorExp(op, def->getOperand(0));

mlir/lib/Dialect/Tensor/IR/TensorOps.cpp

 // doesn't request otherwise.
 struct FoldStaticZeroPadding : public OpRewritePattern<PadOp> {
   using OpRewritePattern<PadOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(PadOp padTensorOp,
                                 PatternRewriter &rewriter) const override {
     if (!padTensorOp.hasZeroLowPad() || !padTensorOp.hasZeroHighPad())
       return failure();
-    if (padTensorOp.getNofold())
+    if (padTensorOp.isNofold())
       return failure();
     rewriter.replaceOpWithNewOp<tensor::CastOp>(
         padTensorOp, padTensorOp.getResult().getType(),
         padTensorOp.getSource());
     return success();
   }
 };
 
LogicalResult matchAndRewrite(PadOp padTensorOp,
     if (newResultType == padTensorOp.getResultType()) {
       rewriter.updateRootInPlace(padTensorOp, [&]() {
         padTensorOp.getSourceMutable().assign(castOp.getSource());
       });
     } else {
       auto newOp = rewriter.create<PadOp>(
           padTensorOp->getLoc(), newResultType, padTensorOp.getSource(),
           padTensorOp.getStaticLow(), padTensorOp.getStaticHigh(),
-          padTensorOp.getLow(), padTensorOp.getHigh(), padTensorOp.getNofold(),
+          padTensorOp.getLow(), padTensorOp.getHigh(), padTensorOp.isNofold(),
           getPrunedAttributeList(padTensorOp, PadOp::getAttributeNames()));
       IRMapping mapper;
       padTensorOp.getRegion().cloneInto(&newOp.getRegion(), mapper);
 
       rewriter.replaceOpWithNewOp<tensor::CastOp>(
           padTensorOp, padTensorOp.getResultType(), newOp);
     }
     return success();
LogicalResult matchAndRewrite(PadOp padTensorOp,
     if (!tensor::preservesStaticInformation(padTensorOp.getResult().getType(),
                                             tensorCastOp.getDest().getType()))
       return failure();
 
     auto replacementOp = rewriter.create<PadOp>(
         padTensorOp.getLoc(), tensorCastOp.getDest().getType(),
         padTensorOp.getSource(), padTensorOp.getStaticLow(),
         padTensorOp.getStaticHigh(), padTensorOp.getLow(),
-        padTensorOp.getHigh(), padTensorOp.getNofold(),
+        padTensorOp.getHigh(), padTensorOp.isNofold(),
         getPrunedAttributeList(padTensorOp, PadOp::getAttributeNames()));
     replacementOp.getRegion().takeBody(padTensorOp.getRegion());
 
     rewriter.replaceOp(padTensorOp, replacementOp.getResult());
     rewriter.replaceOp(tensorCastOp, replacementOp.getResult());
     return success();
   }
 };
struct FoldOrthogonalPaddings : public OpRewritePattern<PadOp> {
   using OpRewritePattern<PadOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(PadOp padOp,
                                 PatternRewriter &rewriter) const override {
     auto innerSliceOp = padOp.getSource().getDefiningOp<ExtractSliceOp>();
     if (!innerSliceOp)
       return failure();
     auto outerPadOp = innerSliceOp.getSource().getDefiningOp<PadOp>();
-    if (!outerPadOp || outerPadOp.getNofold())
+    if (!outerPadOp || outerPadOp.isNofold())
       return failure();
     auto outerSliceOp = outerPadOp.getSource().getDefiningOp<ExtractSliceOp>();
     if (!outerSliceOp)
       return failure();
 
     // 1) Fail if the chain is rank-reducing.
     int64_t rank = padOp.getSourceType().getRank();
     if (outerSliceOp.getSourceType().getRank() != rank) {
LogicalResult matchAndRewrite(PadOp padOp,
 
     // Create a new tensor::ExtractSliceOp, tensor::PadOp pair that performs
     // the two paddings in one step.
     auto newSliceOp = rewriter.create<ExtractSliceOp>(
         padOp.getLoc(), outerSliceOp.getSource(), newOffsets, newSizes,
         innerSliceOp.getMixedStrides());
     auto newPadOp = rewriter.create<PadOp>(
         padOp.getLoc(), padOp.getResultType(), newSliceOp.getResult(),
-        padOp.getMixedLowPad(), newHighPad, padOp.getNofold(),
+        padOp.getMixedLowPad(), newHighPad, padOp.isNofold(),
         getPrunedAttributeList(padOp, PadOp::getAttributeNames()));
     rewriter.inlineRegionBefore(padOp.getRegion(), newPadOp.getRegion(),
                                 newPadOp.getRegion().begin());
     rewriter.replaceOp(padOp, newPadOp.getResult());
     return success();
   }
 };
 
if (SmallVector<int64_t>(outputDims) == newOutDims ||
                      [&](int64_t x) { return x == ShapedType::kDynamic; }))
       return failure();
 
     // Rewrite the op using the new static type.
     auto newResultType = RankedTensorType::get(
         newOutDims, padTensorOp.getType().getElementType());
     auto newOp = rewriter.create<PadOp>(
         padTensorOp->getLoc(), newResultType, input, staticLow, staticHigh,
-        padTensorOp.getLow(), padTensorOp.getHigh(), padTensorOp.getNofold(),
+        padTensorOp.getLow(), padTensorOp.getHigh(), padTensorOp.isNofold(),
         getPrunedAttributeList(padTensorOp, PadOp::getAttributeNames()));
 
     IRMapping mapper;
     padTensorOp.getRegion().cloneInto(&newOp.getRegion(), mapper);
     rewriter.replaceOpWithNewOp<tensor::CastOp>(padTensorOp, oldResultType,
                                                 newOp);
 
     return success();
Value PadOp::getConstantPaddingValue() {
   if (padValue.getParentBlock() == &getRegion().front())
     return {};
   // Else: Yield value defined outside of the PadOp block.
   return padValue;
 }
 
 OpFoldResult PadOp::fold(FoldAdaptor) {
   if (getResultType().hasStaticShape() && getResultType() == getSourceType() &&
-      !getNofold())
+      !isNofold())
     return getSource();
   return {};
 }
 
 //===----------------------------------------------------------------------===//
 // ParallelInsertSliceOp
 //===----------------------------------------------------------------------===//
 

 
 LogicalResult ScatterOp::verify() {
   int64_t destRank = getDestType().getRank();
   ArrayRef<int64_t> scatterDims = getScatterDims();
   if (failed(verifyGatherOrScatterDims(getOperation(), scatterDims, destRank,
                                        "scatter", "dest")))
     return failure();
 
-  if (!getUnique())
+  if (!isUnique())
     return emitOpError("requires 'unique' attribute to be set");
   // TODO: we could also check statically that there are fewer leading index
   // tensor dims than the dest dims. If this is not the case, the unique
   // attribute cannot be true.
 
   // Use the GatherOp::inferResultType on the `dest` type and verify the
   // expected type matches the source type.
   RankedTensorType expectedSourceType = GatherOp::inferResultType(

mlir/lib/Dialect/Tensor/IR/TensorTilingInterfaceImpl.cpp

FailureOr<TilingResult> tensor::bubbleUpPadSlice(OpBuilder &b,
   // Emit a SliceOp and a PadOp. Should not be used in cases where
   // the result shape of the new SliceOp has a zero dimension.
   auto createPadOfExtractSlice = [&]() {
     // Create pad(extract_slice(x)).
     Value newSliceOp = b.create<tensor::ExtractSliceOp>(
         loc, padOp.getSource(), newOffsets, newLengths, newStrides);
     auto newPadOp = b.create<PadOp>(
         loc, Type(), newSliceOp, newLows, newHighs,
-        /*nofold=*/padOp.getNofold(),
+        /*nofold=*/padOp.isNofold(),
         getPrunedAttributeList(padOp, PadOp::getAttributeNames()));
 
     // Copy region to new PadOp.
     IRMapping bvm;
     padOp.getRegion().cloneInto(&newPadOp.getRegion(), bvm);
 
     // Cast result and return.
     return newPadOp;

mlir/lib/Dialect/Tensor/Transforms/FoldIntoPackAndUnpackPatterns.cpp

 /// the pad op has zero low paddings, or if `pack` has no padding values.
 struct FoldPadWithPackOp : public OpRewritePattern<PackOp> {
   using OpRewritePattern<PackOp>::OpRewritePattern;
 
   LogicalResult matchAndRewrite(PackOp packOp,
                                 PatternRewriter &rewriter) const override {
     auto padOp = packOp.getSource().getDefiningOp<PadOp>();
 
-    if (!padOp || padOp.getNofold() || !padOp.hasZeroLowPad())
+    if (!padOp || padOp.isNofold() || !padOp.hasZeroLowPad())
       return failure();
 
     Value constantPaddingValue = padOp.getConstantPaddingValue();
     if (!constantPaddingValue)
       return failure();
 
     if (auto paddingValue = packOp.getPaddingValue())
       if (!isEqualConstantIntOrValue(paddingValue, constantPaddingValue))

mlir/lib/Dialect/Tensor/Transforms/IndependenceTransforms.cpp

FailureOr<Value> tensor::buildIndependentOp(OpBuilder &b, tensor::PadOp padOp,
   // Return existing tensor::PadOp if nothing has changed.
   if (llvm::equal(padOp.getMixedLowPad(), newMixedLow) &&
       llvm::equal(padOp.getMixedHighPad(), newMixedHigh))
     return padOp.getResult();
 
   // Create a new tensor::PadOp.
   auto newPadOp = b.create<PadOp>(
       loc, padOp.getResultType(), padOp.getSource(), newMixedLow, newMixedHigh,
-      constantPadding, padOp.getNofold(), /*attrs=*/ArrayRef<NamedAttribute>{});
+      constantPadding, padOp.isNofold(), /*attrs=*/ArrayRef<NamedAttribute>{});
 
   // Create a tensor::ExtractSliceOp.
   // Reify the result sizes of the old tensor::PadOp.
   ReifiedRankedShapedTypeDims reifiedSizes;
   ReifyRankedShapedTypeOpInterface reifyShapedTypeInterface =
       dyn_cast<ReifyRankedShapedTypeOpInterface>(padOp.getOperation());
   if (failed(reifyShapedTypeInterface.reifyResultShapes(b, reifiedSizes)))
     return failure();

mlir/lib/Dialect/Transform/IR/TransformOps.cpp

 //===----------------------------------------------------------------------===//
 
 DiagnosedSilenceableFailure
 transform::MergeHandlesOp::apply(transform::TransformResults &results,
                                  transform::TransformState &state) {
   SmallVector<Operation *> operations;
   for (Value operand : getHandles())
     llvm::append_range(operations, state.getPayloadOps(operand));
-  if (!getDeduplicate()) {
+  if (!isDeduplicate()) {
     results.set(llvm::cast<OpResult>(getResult()), operations);
     return DiagnosedSilenceableFailure::success();
   }
 
   SetVector<Operation *> uniqued(operations.begin(), operations.end());
   results.set(llvm::cast<OpResult>(getResult()), uniqued.getArrayRef());
   return DiagnosedSilenceableFailure::success();
 }
 
 bool transform::MergeHandlesOp::allowsRepeatedHandleOperands() {
   // Handles may be the same if deduplicating is enabled.
-  return getDeduplicate();
+  return isDeduplicate();
 }
 
 void transform::MergeHandlesOp::getEffects(
     SmallVectorImpl<MemoryEffects::EffectInstance> &effects) {
   onlyReadsHandle(getHandles(), effects);
   producesHandle(getResult(), effects);
 
   // There are no effects on the Payload IR as this is only a handle
   // manipulation.
 }
 
 OpFoldResult transform::MergeHandlesOp::fold(FoldAdaptor adaptor) {
-  if (getDeduplicate() || getHandles().size() != 1)
+  if (isDeduplicate() || getHandles().size() != 1)
     return {};
 
   // If deduplication is not required and there is only one operand, it can be
   // used directly instead of merging.
   return getHandles().front();
 }
 
 //===----------------------------------------------------------------------===//

mlir/lib/Rewrite/ByteCode.cpp

 void Generator::generate(pdl_interp::CheckAttributeOp op,
                          ByteCodeWriter &writer) {
   writer.append(OpCode::AreEqual, op.getAttribute(), op.getConstantValue(),
                 op.getSuccessors());
 }
 void Generator::generate(pdl_interp::CheckOperandCountOp op,
                          ByteCodeWriter &writer) {
   writer.append(OpCode::CheckOperandCount, op.getInputOp(), op.getCount(),
-                static_cast<ByteCodeField>(op.getCompareAtLeast()),
+                static_cast<ByteCodeField>(op.isCompareAtLeast()),
                 op.getSuccessors());
 }
 void Generator::generate(pdl_interp::CheckOperationNameOp op,
                          ByteCodeWriter &writer) {
   writer.append(OpCode::CheckOperationName, op.getInputOp(),
                 OperationName(op.getName(), ctx), op.getSuccessors());
 }
 void Generator::generate(pdl_interp::CheckResultCountOp op,
                          ByteCodeWriter &writer) {
   writer.append(OpCode::CheckResultCount, op.getInputOp(), op.getCount(),
-                static_cast<ByteCodeField>(op.getCompareAtLeast()),
+                static_cast<ByteCodeField>(op.isCompareAtLeast()),
                 op.getSuccessors());
 }
 void Generator::generate(pdl_interp::CheckTypeOp op, ByteCodeWriter &writer) {
   writer.append(OpCode::AreEqual, op.getValue(), op.getType(),
                 op.getSuccessors());
 }
 void Generator::generate(pdl_interp::CheckTypesOp op, ByteCodeWriter &writer) {
   writer.append(OpCode::CheckTypes, op.getValue(), op.getTypes(),
void Generator::generate(pdl_interp::CreateOperationOp op,
   // Add the attributes.
   OperandRange attributes = op.getInputAttributes();
   writer.append(static_cast<ByteCodeField>(attributes.size()));
   for (auto it : llvm::zip(op.getInputAttributeNames(), attributes))
     writer.append(std::get<0>(it), std::get<1>(it));
 
   // Add the result types. If the operation has inferred results, we use a
   // marker "size" value. Otherwise, we add the list of explicit result types.
-  if (op.getInferredResultTypes())
+  if (op.isInferredResultTypes())
     writer.append(kInferTypesMarker);
   else
     writer.appendPDLValueList(op.getInputResultTypes());
 }
 void Generator::generate(pdl_interp::CreateRangeOp op, ByteCodeWriter &writer) {
   // Append the correct opcode for the range type.
   TypeSwitch<Type>(op.getType().getElementType())
       .Case(

mlir/lib/TableGen/Operator.cpp

auto Operator::getArgToOperandOrAttribute(int index) const
     -> OperandOrAttribute {
   return attrOrOperandMapping[index];
 }
 
 std::string Operator::getGetterName(StringRef name) const {
   return "get" + convertToCamelFromSnakeCase(name, /*capitalizeFirst=*/true);
 }
 
+std::string Operator::getGetterNameUnitAttr(StringRef name) const {
+  std::string partialName =
+      convertToCamelFromSnakeCase(name,
+                                  /*capitalizeFirst=*/true);
+  for (auto &namedAttr : getAttributes()) {
+    if (namedAttr.name == name &&
+        namedAttr.attr.getBaseAttr().getAttrDefName() == "UnitAttr") {
+      return "is" + partialName;
+    }
+  }
+  return "get" + partialName;
+}
+
+std::string Operator::getGetterNameAttr(const NamedAttribute &attr) const {
+  std::string name = convertToCamelFromSnakeCase(attr.name,
+                                                 /*capitalizeFirst=*/true);
+  return attr.attr.getBaseAttr().getAttrDefName() == "UnitAttr" ? "is" + name
+                                                                : "get" + name;
+}
+
 std::string Operator::getSetterName(StringRef name) const {
   return "set" + convertToCamelFromSnakeCase(name, /*capitalizeFirst=*/true);
 }
 
 std::string Operator::getRemoverName(StringRef name) const {
   return "remove" + convertToCamelFromSnakeCase(name, /*capitalizeFirst=*/true);
 }
 
 bool Operator::hasFolder() const { return def.getValueAsBit("hasFolder"); }

mlir/lib/Target/Cpp/TranslateToCpp.cpp

static LogicalResult printOperation(CppEmitter &emitter, emitc::CastOp castOp) {
   return success();
 }
 
 static LogicalResult printOperation(CppEmitter &emitter,
                                     emitc::IncludeOp includeOp) {
   raw_ostream &os = emitter.ostream();
 
   os << "#include ";
-  if (includeOp.getIsStandardInclude())
+  if (includeOp.isStandardInclude())
     os << "<" << includeOp.getInclude() << ">";
   else
     os << "\"" << includeOp.getInclude() << "\"";
 
   return success();
 }
 
 static LogicalResult printOperation(CppEmitter &emitter, scf::ForOp forOp) {

mlir/lib/Target/LLVMIR/Dialect/LLVMIR/LLVMToLLVMIRTranslation.cpp

if (auto inlineAsmOp = dyn_cast<LLVM::InlineAsmOp>(opInst)) {
     }
     auto ft = LLVM::LLVMFunctionType::get(resultType, operandTypes);
     llvm::InlineAsm *inlineAsmInst =
         inlineAsmOp.getAsmDialect()
             ? llvm::InlineAsm::get(
                   static_cast<llvm::FunctionType *>(
                       moduleTranslation.convertType(ft)),
                   inlineAsmOp.getAsmString(), inlineAsmOp.getConstraints(),
-                  inlineAsmOp.getHasSideEffects(),
-                  inlineAsmOp.getIsAlignStack(),
+                  inlineAsmOp.isHasSideEffects(), inlineAsmOp.isAlignStack(),
                   convertAsmDialectToLLVM(*inlineAsmOp.getAsmDialect()))
-            : llvm::InlineAsm::get(static_cast<llvm::FunctionType *>(
-                                       moduleTranslation.convertType(ft)),
-                                   inlineAsmOp.getAsmString(),
-                                   inlineAsmOp.getConstraints(),
-                                   inlineAsmOp.getHasSideEffects(),
-                                   inlineAsmOp.getIsAlignStack());
+            : llvm::InlineAsm::get(
+                  static_cast<llvm::FunctionType *>(
+                  inlineAsmOp.getAsmString(), inlineAsmOp.getConstraints(),
+                  inlineAsmOp.isHasSideEffects(), inlineAsmOp.isAlignStack());
     llvm::CallInst *inst = builder.CreateCall(
         inlineAsmInst,
         moduleTranslation.lookupValues(inlineAsmOp.getOperands()));
     if (auto maybeOperandAttrs = inlineAsmOp.getOperandAttrs()) {
       llvm::AttributeList attrList;
       for (const auto &it : llvm::enumerate(*maybeOperandAttrs)) {
         Attribute attr = it.value();
         if (!attr)
if (auto invOp = dyn_cast<LLVM::InvokeOp>(opInst)) {
     }
     return success(result->getType()->isVoidTy());
   }
 
   if (auto lpOp = dyn_cast<LLVM::LandingpadOp>(opInst)) {
     llvm::Type *ty = moduleTranslation.convertType(lpOp.getType());
     llvm::LandingPadInst *lpi =
         builder.CreateLandingPad(ty, lpOp.getNumOperands());
-    lpi->setCleanup(lpOp.getCleanup());
+    lpi->setCleanup(lpOp.isCleanup());
 
     // Add clauses
     for (llvm::Value *operand :
          moduleTranslation.lookupValues(lpOp.getOperands())) {
       // All operands should be constant - checked by verifier
       if (auto *constOperand = dyn_cast<llvm::Constant>(operand))
         lpi->addClause(constOperand);
     }

mlir/lib/Target/LLVMIR/Dialect/OpenMP/OpenMPToLLVMIRTranslation.cpp

 /// OpenMPIRBuilder.
 static LogicalResult
 convertOmpOrderedRegion(Operation &opInst, llvm::IRBuilderBase &builder,
                         LLVM::ModuleTranslation &moduleTranslation) {
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
   auto orderedRegionOp = cast<omp::OrderedRegionOp>(opInst);
 
   // TODO: The code generation for ordered simd directive is not supported yet.
-  if (orderedRegionOp.getSimd())
+  if (orderedRegionOp.isSimd())
     return failure();
 
   // TODO: support error propagation in OpenMPIRBuilder and use it instead of
   // relying on captured variables.
   LogicalResult bodyGenStatus = success();
 
   auto bodyGenCB = [&](InsertPointTy allocaIP, InsertPointTy codeGenIP) {
     // OrderedOp has only one region associated with it.
     auto &region = cast<omp::OrderedRegionOp>(opInst).getRegion();
     builder.restoreIP(codeGenIP);
     convertOmpOpRegions(region, "omp.ordered.region", builder,
                         moduleTranslation, bodyGenStatus);
   };
 
   // TODO: Perform finalization actions for variables. This has to be
   // called for variables which have destructors/finalizers.
   auto finiCB = [&](InsertPointTy codeGenIP) {};
 
   llvm::OpenMPIRBuilder::LocationDescription ompLoc(builder);
   builder.restoreIP(
       moduleTranslation.getOpenMPBuilder()->createOrderedThreadsSimd(
-          ompLoc, bodyGenCB, finiCB, !orderedRegionOp.getSimd()));
+          ompLoc, bodyGenCB, finiCB, !orderedRegionOp.isSimd()));
   return bodyGenStatus;
 }
 
 static LogicalResult
 convertOmpSections(Operation &opInst, llvm::IRBuilderBase &builder,
                    LLVM::ModuleTranslation &moduleTranslation) {
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
   using StorableBodyGenCallbackTy =
convertOmpSections(Operation &opInst, llvm::IRBuilderBase &builder,
   // called for variables which have destructors/finalizers.
   auto finiCB = [&](InsertPointTy codeGenIP) {};
 
   llvm::OpenMPIRBuilder::InsertPointTy allocaIP =
       findAllocaInsertPoint(builder, moduleTranslation);
   llvm::OpenMPIRBuilder::LocationDescription ompLoc(builder);
   builder.restoreIP(moduleTranslation.getOpenMPBuilder()->createSections(
       ompLoc, allocaIP, sectionCBs, privCB, finiCB, false,
-      sectionsOp.getNowait()));
+      sectionsOp.isNowait()));
   return bodyGenStatus;
 }
 
 /// Converts an OpenMP single construct into LLVM IR using OpenMPIRBuilder.
 static LogicalResult
 convertOmpSingle(omp::SingleOp &singleOp, llvm::IRBuilderBase &builder,
                  LLVM::ModuleTranslation &moduleTranslation) {
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
   llvm::OpenMPIRBuilder::LocationDescription ompLoc(builder);
   LogicalResult bodyGenStatus = success();
   auto bodyCB = [&](InsertPointTy allocaIP, InsertPointTy codegenIP) {
     builder.restoreIP(codegenIP);
     convertOmpOpRegions(singleOp.getRegion(), "omp.single.region", builder,
                         moduleTranslation, bodyGenStatus);
   };
   auto finiCB = [&](InsertPointTy codeGenIP) {};
   builder.restoreIP(moduleTranslation.getOpenMPBuilder()->createSingle(
-      ompLoc, bodyCB, finiCB, singleOp.getNowait(), /*DidIt=*/nullptr));
+      ompLoc, bodyCB, finiCB, singleOp.isNowait(), /*DidIt=*/nullptr));
   return bodyGenStatus;
 }
 
 /// Converts an OpenMP task construct into LLVM IR using OpenMPIRBuilder.
 static LogicalResult
 convertOmpTaskOp(omp::TaskOp taskOp, llvm::IRBuilderBase &builder,
                  LLVM::ModuleTranslation &moduleTranslation) {
   using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
   LogicalResult bodyGenStatus = success();
-  if (taskOp.getIfExpr() || taskOp.getFinalExpr() || taskOp.getUntiedAttr() ||
-      taskOp.getMergeableAttr() || taskOp.getInReductions() ||
+  if (taskOp.getIfExpr() || taskOp.getFinalExpr() || taskOp.isUntiedAttr() ||
+      taskOp.isMergeableAttr() || taskOp.getInReductions() ||
       taskOp.getPriority() || !taskOp.getAllocateVars().empty()) {
     return taskOp.emitError("unhandled clauses for translation to LLVM IR");
   }
   auto bodyCB = [&](InsertPointTy allocaIP, InsertPointTy codegenIP) {
     builder.restoreIP(codegenIP);
     convertOmpOpRegions(taskOp.getRegion(), "omp.task.region", builder,
                         moduleTranslation, bodyGenStatus);
   };
for (auto dep :
       dds.emplace_back(dd);
     }
   }
 
   llvm::OpenMPIRBuilder::InsertPointTy allocaIP =
       findAllocaInsertPoint(builder, moduleTranslation);
   llvm::OpenMPIRBuilder::LocationDescription ompLoc(builder);
   builder.restoreIP(moduleTranslation.getOpenMPBuilder()->createTask(
-      ompLoc, allocaIP, bodyCB, !taskOp.getUntied(), /*Final*/ nullptr,
+      ompLoc, allocaIP, bodyCB, !taskOp.isUntied(), /*Final*/ nullptr,
       /*IfCondition*/ nullptr, dds));
   return bodyGenStatus;
 }
 
 /// Converts an OpenMP taskgroup construct into LLVM IR using OpenMPIRBuilder.
 static LogicalResult
 convertOmpTaskgroupOp(omp::TaskGroupOp tgOp, llvm::IRBuilderBase &builder,
                       LLVM::ModuleTranslation &moduleTranslation) {
for (unsigned i = 0, e = loop.getNumLoops(); i < e; ++i) {
     llvm::OpenMPIRBuilder::LocationDescription loc = ompLoc;
     llvm::OpenMPIRBuilder::InsertPointTy computeIP = ompLoc.IP;
     if (i != 0) {
       loc = llvm::OpenMPIRBuilder::LocationDescription(bodyInsertPoints.back());
       computeIP = loopInfos.front()->getPreheaderIP();
     }
     loopInfos.push_back(ompBuilder->createCanonicalLoop(
         loc, bodyGen, lowerBound, upperBound, step,
-        /*IsSigned=*/true, loop.getInclusive(), computeIP));
+        /*IsSigned=*/true, loop.isInclusive(), computeIP));
 
     if (failed(bodyGenStatus))
       return failure();
   }
 
   // Collapse loops. Store the insertion point because LoopInfos may get
   // invalidated.
   llvm::IRBuilderBase::InsertPoint afterIP = loopInfos.front()->getAfterIP();
   llvm::CanonicalLoopInfo *loopInfo =
       ompBuilder->collapseLoops(ompLoc.DL, loopInfos, {});
 
   allocaIP = findAllocaInsertPoint(builder, moduleTranslation);
 
   // TODO: Handle doacross loops when the ordered clause has a parameter.
   bool isOrdered = loop.getOrderedVal().has_value();
   std::optional<omp::ScheduleModifier> scheduleModifier =
       loop.getScheduleModifier();
-  bool isSimd = loop.getSimdModifier();
+  bool isSimd = loop.isSimdModifier();
 
   ompBuilder->applyWorkshareLoop(
-      ompLoc.DL, loopInfo, allocaIP, !loop.getNowait(),
+      ompLoc.DL, loopInfo, allocaIP, !loop.isNowait(),
       convertToScheduleKind(schedule), chunk, isSimd,
       scheduleModifier == omp::ScheduleModifier::monotonic,
       scheduleModifier == omp::ScheduleModifier::nonmonotonic, isOrdered);
 
   // Continue building IR after the loop. Note that the LoopInfo returned by
   // `collapseLoops` points inside the outermost loop and is intended for
   // potential further loop transformations. Use the insertion point stored
   // before collapsing loops instead.
convertOmpWsLoop(Operation &opInst, llvm::IRBuilderBase &builder,
 
   // The call to createReductions below expects the block to have a
   // terminator. Create an unreachable instruction to serve as terminator
   // and remove it later.
   llvm::UnreachableInst *tempTerminator = builder.CreateUnreachable();
   builder.SetInsertPoint(tempTerminator);
   llvm::OpenMPIRBuilder::InsertPointTy contInsertPoint =
       ompBuilder->createReductions(builder.saveIP(), allocaIP, reductionInfos,
-                                   loop.getNowait());
+                                   loop.isNowait());
   if (!contInsertPoint.getBlock())
     return loop->emitOpError() << "failed to convert reductions";
   auto nextInsertionPoint =
       ompBuilder->createBarrier(contInsertPoint, llvm::omp::OMPD_for);
   tempTerminator->eraseFromParent();
   builder.restoreIP(nextInsertionPoint);
 
   return success();
LogicalResult result =
                   deviceID = intAttr.getInt();
 
             numMapOperands = dataOp.getMapOperands().size();
             mapOperands = dataOp.getMapOperands();
             mapTypes = dataOp.getMapTypes();
             return success();
           })
           .Case([&](omp::EnterDataOp enterDataOp) {
-            if (enterDataOp.getNowait())
+            if (enterDataOp.isNowait())
               return failure();
 
             if (auto ifExprVar = enterDataOp.getIfExpr())
               ifCond = moduleTranslation.lookupValue(ifExprVar);
 
             if (auto devId = enterDataOp.getDevice())
               if (auto constOp = mlir::dyn_cast<mlir::LLVM::ConstantOp>(
                       devId.getDefiningOp()))
                 if (auto intAttr =
                         dyn_cast<mlir::IntegerAttr>(constOp.getValue()))
                   deviceID = intAttr.getInt();
 
             numMapOperands = enterDataOp.getMapOperands().size();
             mapOperands = enterDataOp.getMapOperands();
             mapTypes = enterDataOp.getMapTypes();
             return success();
           })
           .Case([&](omp::ExitDataOp exitDataOp) {
-            if (exitDataOp.getNowait())
+            if (exitDataOp.isNowait())
               return failure();
 
             if (auto ifExprVar = exitDataOp.getIfExpr())
               ifCond = moduleTranslation.lookupValue(ifExprVar);
 
             if (auto devId = exitDataOp.getDevice())
               if (auto constOp = mlir::dyn_cast<mlir::LLVM::ConstantOp>(
                       devId.getDefiningOp()))
if (targetOp.getDevice()) {
     return false;
   }
 
   if (targetOp.getThreadLimit()) {
     opInst.emitError("Thread limit clause not yet supported");
     return false;
   }
 
-  if (targetOp.getNowait()) {
+  if (targetOp.isNowait()) {
     opInst.emitError("Nowait clause not yet supported");
     return false;
   }
 
   return true;
 }
 
 static LogicalResult

mlir/lib/Target/LLVMIR/ModuleTranslation.cpp

for (auto op : getModuleBody(mlirModule).getOps<LLVM::GlobalOp>()) {
     // default to undef.
     bool dropInitializer = shouldDropGlobalInitializer(linkage, cst);
     if (!dropInitializer && !cst)
       cst = llvm::UndefValue::get(type);
     else if (dropInitializer && cst)
       cst = nullptr;
 
     auto *var = new llvm::GlobalVariable(
-        *llvmModule, type, op.getConstant(), linkage, cst, op.getSymName(),
+        *llvmModule, type, op.isConstant(), linkage, cst, op.getSymName(),
         /*InsertBefore=*/nullptr,
-        op.getThreadLocal_() ? llvm::GlobalValue::GeneralDynamicTLSModel
-                             : llvm::GlobalValue::NotThreadLocal,
+        op.isThreadLocal_() ? llvm::GlobalValue::GeneralDynamicTLSModel
         addrSpace);
 
     if (op.getUnnamedAddr().has_value())
       var->setUnnamedAddr(convertUnnamedAddrToLLVM(*op.getUnnamedAddr()));
 
     if (op.getSection().has_value())
       var->setSection(*op.getSection());
 
-    addRuntimePreemptionSpecifier(op.getDsoLocal(), var);
+    addRuntimePreemptionSpecifier(op.isDsoLocal(), var);
 
     std::optional<uint64_t> alignment = op.getAlignment();
     if (alignment.has_value())
       var->setAlignment(llvm::MaybeAlign(alignment.value()));
 
     var->setVisibility(convertVisibilityToLLVM(op.getVisibility_()));
 
     globalsMapping.try_emplace(op, var);
LogicalResult ModuleTranslation::convertFunctionSignatures() {
   for (auto function : getModuleBody(mlirModule).getOps<LLVMFuncOp>()) {
     llvm::FunctionCallee llvmFuncCst = llvmModule->getOrInsertFunction(
         function.getName(),
         cast<llvm::FunctionType>(convertType(function.getFunctionType())));
     llvm::Function *llvmFunc = cast<llvm::Function>(llvmFuncCst.getCallee());
     llvmFunc->setLinkage(convertLinkageToLLVM(function.getLinkage()));
     llvmFunc->setCallingConv(convertCConvToLLVM(function.getCConv()));
     mapFunction(function.getName(), llvmFunc);
-    addRuntimePreemptionSpecifier(function.getDsoLocal(), llvmFunc);
+    addRuntimePreemptionSpecifier(function.isDsoLocal(), llvmFunc);
 
     // Convert function attributes.
     convertFunctionAttributes(function, llvmFunc);
 
     // Convert function_entry_count attribute to metadata.
     if (std::optional<uint64_t> entryCount = function.getFunctionEntryCount())
       llvmFunc->setEntryCount(entryCount.value());
 
for (auto &op : metadata.getBody().getOps()) {
         // Use temporary MDNode's as the placeholders for the base and access
         // types to prevent uniquing the tag nodes until they are finalized.
         operands.push_back(
             llvm::MDNode::getTemporary(ctx, std::nullopt).release());
         operands.push_back(
             llvm::MDNode::getTemporary(ctx, std::nullopt).release());
         operands.push_back(llvm::ConstantAsMetadata::get(
             llvm::ConstantInt::get(offsetTy, tagOp.getOffset())));
-        if (tagOp.getConstant())
+        if (tagOp.isConstant())
           operands.push_back(llvm::ConstantAsMetadata::get(
               llvm::ConstantInt::get(offsetTy, 1)));
       }
 
       if (operands.empty())
         continue;
 
       tbaaMetadataMapping.insert({&op, llvm::MDNode::get(ctx, operands)});

mlir/test/Dialect/LLVMIR/roundtrip.mlir

llvm.func @useInlineAsm(%arg0: i32) {
   %0 = llvm.inline_asm "bswap $0", "=r,r" %arg0 : (i32) -> i8
 
   // CHECK-NEXT:  llvm.inline_asm {{.*}} (i32, i32) -> i8
   %1 = llvm.inline_asm "foo", "bar" %arg0, %arg0 : (i32, i32) -> i8
 
   // CHECK-NEXT:  llvm.inline_asm has_side_effects {{.*}} (i32, i32) -> i8
   %2 = llvm.inline_asm has_side_effects "foo", "bar" %arg0, %arg0 : (i32, i32) -> i8
 
-  // CHECK-NEXT:  llvm.inline_asm is_align_stack {{.*}} (i32, i32) -> i8
-  %3 = llvm.inline_asm is_align_stack "foo", "bar" %arg0, %arg0 : (i32, i32) -> i8
+  // CHECK-NEXT:  llvm.inline_asm align_stack {{.*}} (i32, i32) -> i8
+  %3 = llvm.inline_asm align_stack "foo", "bar" %arg0, %arg0 : (i32, i32) -> i8
 
   // CHECK-NEXT:  llvm.inline_asm "foo", "=r,=r,r" {{.*}} : (i32) -> !llvm.struct<(i8, i8)>
   %5 = llvm.inline_asm "foo", "=r,=r,r" %arg0 : (i32) -> !llvm.struct<(i8, i8)>
 
   llvm.return
 }
 
 // CHECK-LABEL: @fastmathFlags

mlir/test/Dialect/Linalg/match-ops-invalid.mlir

 
 // -----
 
 transform.sequence failures(suppress) {
 ^bb0(%arg0: !transform.any_op):
   transform.match.structured %arg0 : !transform.any_op {
   ^bb1(%arg1: !transform.any_op):
     // expected-error @below {{cannot request both 'all' and 'inverted' values in the list}}
-    "transform.match.structured.dim"(%arg1) { is_all, is_inverted, raw_dim_list = array<i64> } : (!transform.any_op) -> ()
+    "transform.match.structured.dim"(%arg1) { all, inverted, raw_dim_list = array<i64> } : (!transform.any_op) -> ()
     transform.match.structured.yield
   }
   transform.yield
 }
 
 // -----
 
 transform.sequence failures(suppress) {
 ^bb0(%arg0: !transform.any_op):
   transform.match.structured %arg0 : !transform.any_op {
   ^bb1(%arg1: !transform.any_op):
     // expected-error @below {{cannot both request 'all' and specific values in the list}}
-    "transform.match.structured.dim"(%arg1) { is_all, raw_dim_list = array<i64: 0, 1> } : (!transform.any_op) -> ()
+    "transform.match.structured.dim"(%arg1) { all, raw_dim_list = array<i64: 0, 1> } : (!transform.any_op) -> ()
     transform.match.structured.yield
   }
   transform.yield
 }
 // -----
 
 transform.sequence failures(suppress) {
 ^bb0(%arg0: !transform.any_op):

 
 // -----
 
 transform.sequence failures(suppress) {
 ^bb0(%arg0: !transform.any_op):
   transform.match.structured %arg0 : !transform.any_op {
   ^bb1(%arg1: !transform.any_op):   
     // expected-error @below {{cannot request the same dimension to be both parallel and reduction}}
-    "transform.match.structured.dim"(%arg1) { is_all, parallel, reduction, raw_dim_list = array<i64> } : (!transform.any_op) -> ()
+    "transform.match.structured.dim"(%arg1) { all, parallel, reduction, raw_dim_list = array<i64> } : (!transform.any_op) -> ()
     transform.match.structured.yield
   }
   transform.yield
 }
 
 // -----
 
 transform.sequence failures(suppress) {

mlir/test/Target/LLVMIR/llvmir.mlir

llvm.func @useInlineAsm(%arg0: i32) {
 
   // CHECK-NEXT:  call i8 asm "foo", "=r,r,r"(i32 {{.*}}, i32 {{.*}})
   %1 = llvm.inline_asm "foo", "=r,r,r" %arg0, %arg0 : (i32, i32) -> i8
 
   // CHECK-NEXT:  call i8 asm sideeffect "foo", "=r,r,r"(i32 {{.*}}, i32 {{.*}})
   %2 = llvm.inline_asm has_side_effects "foo", "=r,r,r" %arg0, %arg0 : (i32, i32) -> i8
 
   // CHECK-NEXT:  call i8 asm alignstack "foo", "=r,r,r"(i32 {{.*}}, i32 {{.*}})
-  %3 = llvm.inline_asm is_align_stack "foo", "=r,r,r" %arg0, %arg0 : (i32, i32) -> i8
+  %3 = llvm.inline_asm align_stack "foo", "=r,r,r" %arg0, %arg0 : (i32, i32) -> i8
 
   // CHECK-NEXT:  call i8 asm inteldialect "foo", "=r,r,r"(i32 {{.*}}, i32 {{.*}})
   %4 = llvm.inline_asm asm_dialect = "intel" "foo", "=r,r,r" %arg0, %arg0 : (i32, i32) -> i8
 
   // CHECK-NEXT:  call { i8, i8 } asm "foo", "=r,=r,r"(i32 {{.*}})
   %5 = llvm.inline_asm "foo", "=r,=r,r" %arg0 : (i32) -> !llvm.struct<(i8, i8)>
 
   llvm.return

mlir/test/lib/Dialect/Test/TestOps.td

 def FormatOptionalDefaultAttrs : TEST_Op<"format_optional_default_attrs"> {
   let arguments = (ins DefaultValuedStrAttr<StrAttr, "default">:$str,
                        DefaultValuedStrAttr<SymbolNameAttr, "default">:$sym,
                        DefaultValuedAttr<SomeI64Enum, "SomeI64Enum::case5">:$e);
   let assemblyFormat = "($str^)? ($sym^)? ($e^)? attr-dict";
 }
 
 def FormatOptionalWithElse : TEST_Op<"format_optional_else"> {
-  let arguments = (ins UnitAttr:$isFirstBranchPresent);
-  let assemblyFormat = "(`then` $isFirstBranchPresent^):(`else`)? attr-dict";
+  let arguments = (ins UnitAttr:$firstBranchPresent);
+  let assemblyFormat = "(`then` $firstBranchPresent^):(`else`)? attr-dict";
 }
 
 def FormatCompoundAttr : TEST_Op<"format_compound_attr"> {
   let arguments = (ins CompoundAttrA:$compound);
   let assemblyFormat = "$compound attr-dict-with-keyword";
 }
 
 def FormatNestedAttr : TEST_Op<"format_nested_attr"> {

mlir/test/lib/Dialect/Transform/TestTransformDialectExtension.cpp

void mlir::test::TestBranchingTransformOpTerminator::getEffects(
     SmallVectorImpl<MemoryEffects::EffectInstance> &effects) {}
 
 DiagnosedSilenceableFailure mlir::test::TestEmitRemarkAndEraseOperandOp::apply(
     transform::TransformResults &results, transform::TransformState &state) {
   emitRemark() << getRemark();
   for (Operation *op : state.getPayloadOps(getTarget()))
     op->erase();
 
-  if (getFailAfterErase())
+  if (isFailAfterErase())
     return emitSilenceableError() << "silenceable error";
   return DiagnosedSilenceableFailure::success();
 }
 
 void mlir::test::TestEmitRemarkAndEraseOperandOp::getEffects(
     SmallVectorImpl<MemoryEffects::EffectInstance> &effects) {
   transform::consumesHandle(getTarget(), effects);
   transform::modifiesPayload(effects);
void mlir::test::TestProduceTransformParamOrForwardOperandOp::getEffects(
   transform::producesHandle(getParam(), effects);
 }
 
 DiagnosedSilenceableFailure
 mlir::test::TestProduceTransformParamOrForwardOperandOp::applyToOne(
     Operation *target, ::transform::ApplyToEachResultList &results,
     ::transform::TransformState &state) {
   Builder builder(getContext());
-  if (getFirstResultIsParam()) {
+  if (isFirstResultIsParam()) {
     results.push_back(builder.getI64IntegerAttr(0));
-  } else if (getFirstResultIsNull()) {
+  } else if (isFirstResultIsNull()) {
     results.push_back(nullptr);
   } else {
     results.push_back(*state.getPayloadOps(getIn()).begin());
   }
 
-  if (getSecondResultIsHandle()) {
+  if (isSecondResultIsHandle()) {
     results.push_back(*state.getPayloadOps(getIn()).begin());
   } else {
     results.push_back(builder.getI64IntegerAttr(42));
   }
 
   return DiagnosedSilenceableFailure::success();
 }
 

 mlir::test::TestProduceNullValueOp::apply(transform::TransformResults &results,
                                           transform::TransformState &state) {
   results.setValues(llvm::cast<OpResult>(getOut()), Value());
   return DiagnosedSilenceableFailure::success();
 }
 
 void mlir::test::TestRequiredMemoryEffectsOp::getEffects(
     SmallVectorImpl<MemoryEffects::EffectInstance> &effects) {
-  if (getHasOperandEffect())
+  if (isHasOperandEffect())
     transform::consumesHandle(getIn(), effects);
 
-  if (getHasResultEffect())
+  if (isHasResultEffect())
     transform::producesHandle(getOut(), effects);
   else
     transform::onlyReadsHandle(getOut(), effects);
 
-  if (getModifiesPayload())
+  if (isModifiesPayload())
     transform::modifiesPayload(effects);
 }
 
 DiagnosedSilenceableFailure mlir::test::TestRequiredMemoryEffectsOp::apply(
     transform::TransformResults &results, transform::TransformState &state) {
   results.set(llvm::cast<OpResult>(getOut()), state.getPayloadOps(getIn()));
   return DiagnosedSilenceableFailure::success();
 }

mlir/test/mlir-tblgen/op-attribute.td

 // Test unit attributes.
 // ---
 
 def UnitAttrOp : NS_Op<"unit_attr_op", []> {
   let arguments = (ins UnitAttr:$attr);
 }
 
 // DEF-LABEL: UnitAttrOp definitions
-// DEF: bool UnitAttrOp::getAttr() {
+// DEF: bool UnitAttrOp::isAttr() {
 // DEF:   return {{.*}} != nullptr
 
 // DEF: ::mlir::Attribute UnitAttrOp::removeAttrAttr() {
 // DEF-NEXT:   (*this)->removeAttr(getAttrAttrName());
 
 // DEF: build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, /*optional*/::mlir::UnitAttr attr)
 // DEF: build(::mlir::OpBuilder &odsBuilder, ::mlir::OperationState &odsState, /*optional*/bool attr)
 

mlir/test/mlir-tblgen/op-format.td

 }]>, Arguments<(ins Optional<I1>:$a)>;
 
 // CHECK-LABEL: OptionalGroupB::parse
 // CHECK: if (::mlir::succeeded(parser.parseOptionalKeyword("foo")))
 // CHECK-NEXT: else
 // CHECK-NEXT: result.addAttribute("a", parser.getBuilder().getUnitAttr())
 // CHECK: parser.parseKeyword("bar")
 // CHECK-LABEL: OptionalGroupB::print
-// CHECK: if (!getAAttr())
+// CHECK: if (!isAAttr())
 // CHECK-NEXT: odsPrinter << ' ' << "foo"
 // CHECK-NEXT: else
 // CHECK-NEXT: odsPrinter << ' ' << "bar"
 def OptionalGroupB : TestFormat_Op<[{
   (`foo`) : (`bar` $a^)? attr-dict
 }]>, Arguments<(ins UnitAttr:$a)>;
 
 // Optional group anchored on a default-valued attribute:

mlir/tools/mlir-tblgen/LLVMIRConversionGen.cpp

static LogicalResult emitOneBuilder(const Record &record, raw_ostream &os) {
 
   // Progressively create the builder string by replacing $-variables with
   // value lookups.  Keep only the not-yet-traversed part of the builder pattern
   // to avoid re-traversing the string multiple times.
   std::string builder;
   llvm::raw_string_ostream bs(builder);
   while (StringLoc loc = findNextVariable(builderStrRef)) {
     auto name = loc.in(builderStrRef).drop_front();
-    auto getterName = op.getGetterName(name);
+    auto getterName = op.getGetterNameUnitAttr(name);
     // First, insert the non-matched part as is.
     bs << builderStrRef.substr(0, loc.pos);
     // Then, rewrite the name based on its kind.
     bool isVariadicOperand = isVariadicOperandName(op, name);
     if (isOperandName(op, name)) {
       auto result =
           isVariadicOperand
               ? formatv("moduleTranslation.lookupValues(op.{0}())", getterName)

mlir/tools/mlir-tblgen/OpDefinitionsGen.cpp

 // Populate the format context `ctx` with substitutions of attributes, operands
 // and results.
 static void populateSubstitutions(const OpOrAdaptorHelper &emitHelper,
                                   FmtContext &ctx) {
   // Populate substitutions for attributes.
   auto &op = emitHelper.getOp();
   for (const auto &namedAttr : op.getAttributes())
     ctx.addSubst(namedAttr.name,
-                 emitHelper.getOp().getGetterName(namedAttr.name) + "()");
+                 emitHelper.getOp().getGetterNameAttr(namedAttr) + "()");
 
   // Populate substitutions for named operands.
   for (int i = 0, e = op.getNumOperands(); i < e; ++i) {
     auto &value = op.getOperand(i);
     if (!value.name.empty())
       ctx.addSubst(value.name, emitHelper.getOperand(i).str());
   }
 
for (const std::pair<StringRef, AttributeMetadata> &it :
         continue;
       }
 
       body << formatv("::mlir::Attribute {0};\n", getVarName(it.first));
       for (const AttributeMetadata *optional : optionalAttrs) {
         body << formatv("::mlir::Attribute {0};\n",
                         getVarName(optional->attrName));
       }
-      body << formatv(findRequiredAttr, emitHelper.getAttrName(it.first),
+      body << formatv(findRequiredAttr,
+                      emitHelper.getAttrName(it.first),
                       emitHelper.emitErrorPrefix(), it.first);
       for (const AttributeMetadata *optional : optionalAttrs) {
         body << formatv(checkOptionalAttr,
                         emitHelper.getAttrName(optional->attrName),
                         optional->attrName);
       }
       body << "\n  ++namedAttrIt;\n}\n";
       optionalAttrs.clear();

 static void emitAttrGetterWithReturnType(FmtContext &fctx,
                                          OpClassOrAdaptor &opClass,
                                          const Operator &op, StringRef name,
                                          Attribute attr) {
   auto *method = opClass.addMethod(attr.getReturnType(), name);
   ERROR_IF_PRUNED(method, name, op);
   auto &body = method->body();
   body << "  auto attr = " << name << "Attr();\n";
   if (attr.hasDefaultValue() && attr.isOptional()) {
     // Returns the default value if not set.
     // TODO: this is inefficient, we are recreating the attribute for every
     // call. This should be set instead.
     if (!attr.isConstBuildable()) {
       PrintFatalError("DefaultValuedAttr of type " + attr.getAttrDefName() +
                       " must have a constBuilder");
     }
     std::string defaultValue = std::string(
         tgfmt(attr.getConstBuilderTemplate(), &fctx, attr.getDefaultValue()));
     body << "    if (!attr)\n      return "
          << tgfmt(attr.getConvertFromStorageCall(),
                   &fctx.withSelf(defaultValue))
          << ";\n";
   }
   body << "  return "
        << tgfmt(attr.getConvertFromStorageCall(), &fctx.withSelf("attr"))
        << ";\n";
 }
 
+template <typename OpClassOrAdaptor>
+static void emitAttrGetterWithReturnTypeDeprecated(
+    FmtContext &fctx, OpClassOrAdaptor &opClass, const Operator &op,
+    StringRef name, StringRef deprecatedName, Attribute attr) {
+  auto *method = opClass.addMethod(attr.getReturnType(), deprecatedName);
+  std::string deprecatedMessage = std::string(deprecatedName) +
+                                  " is deprecated. Use " + std::string(name) +
+                                  " instead.";
+  method->setDeprecated(deprecatedMessage);
+  ERROR_IF_PRUNED(method, deprecatedName, op);
+  auto &body = method->body();
+  body << "  auto attr = " << deprecatedName << "Attr();\n";
+  if (attr.hasDefaultValue() && attr.isOptional()) {
+    // Returns the default value if not set.
+    // TODO: this is inefficient, we are recreating the attribute for every
+    // call. This should be set instead.
+    if (!attr.isConstBuildable()) {
+      PrintFatalError("DefaultValuedAttr of type " + attr.getAttrDefName() +
+                      " must have a constBuilder");
+    }
+    std::string defaultValue = std::string(
+        tgfmt(attr.getConstBuilderTemplate(), &fctx, attr.getDefaultValue()));
+    body << "    if (!attr)\n      return "
+         << tgfmt(attr.getConvertFromStorageCall(),
+                  &fctx.withSelf(defaultValue))
+         << ";\n";
+  }
+  body << "  return "
+       << tgfmt(attr.getConvertFromStorageCall(), &fctx.withSelf("attr"))
+       << ";\n";
+}
+
 void OpEmitter::genPropertiesSupport() {
   if (!emitHelper.hasProperties())
     return;
   using ConstArgument =
       llvm::PointerUnion<const AttributeMetadata *, const NamedProperty *>;
 
   SmallVector<ConstArgument> attrOrProperties;
   for (const std::pair<StringRef, AttributeMetadata> &it :
if (!namedAttr || !namedAttr->constraint)
       continue;
     Attribute attr = *namedAttr->constraint;
     std::optional<StringRef> constraintFn =
         staticVerifierEmitter.getAttrConstraintFn(attr);
     if (!constraintFn)
       continue;
     if (canEmitAttrVerifier(attr,
                             /*isEmittingForOp=*/false)) {
-      std::string name = op.getGetterName(namedAttr->attrName);
+      std::string name =
+          op.getGetterName(namedAttr->attrName);
       verifyInherentAttrsMethod
           << formatv(R"(
     {{
       ::mlir::Attribute attr = attrs.get({0}AttrName(opName));
       if (attr && ::mlir::failed({1}(attr, "{2}", getDiag)))
         return ::mlir::failure();
     }
 )",
                      name, constraintFn, namedAttr->attrName);
     }
   }
   verifyInherentAttrsMethod << "    return ::mlir::success();";
 }
 
 void OpEmitter::genAttrGetters() {
   FmtContext fctx;
   fctx.withBuilder("::mlir::Builder((*this)->getContext())");
 
   // Emit the derived attribute body.
   auto emitDerivedAttr = [&](StringRef name, Attribute attr) {
     if (auto *method = opClass.addMethod(attr.getReturnType(), name))
       method->body() << "  " << attr.getDerivedCodeBody() << "\n";
   };
 
+  // Emit the deprecated derived attribute body for unit attributes
+  // TODO: Method will be removed 2023-06-23/
+  auto emitDerivedAttrDeprecated = [&](StringRef name, StringRef deprecatedName,
+                                       Attribute attr) {
+    if (auto *method =
+            opClass.addMethod(attr.getReturnType(), deprecatedName)) {
+      method->body() << "  " << attr.getDerivedCodeBody() << "\n";
+      std::string deprecatedMessage = std::string(deprecatedName) +
+                                      " is deprecated. Use " +
+                                      std::string(name) + " instead.";
+      method->setDeprecated(deprecatedMessage);
+    }
+  };
+
   // Generate named accessor with Attribute return type. This is a wrapper
   // class that allows referring to the attributes via accessors instead of
   // having to use the string interface for better compile time verification.
   auto emitAttrWithStorageType = [&](StringRef name, StringRef attrName,
                                      Attribute attr) {
     auto *method = opClass.addMethod(attr.getStorageType(), name + "Attr");
     if (!method)
       return;
     method->body() << formatv(
         "  return {0}.{1}<{2}>();", emitHelper.getAttr(attrName),
         attr.isOptional() || attr.hasDefaultValue() ? "dyn_cast_or_null"
                                                     : "cast",
         attr.getStorageType());
   };
 
+  // Generate deprecated named accessor with Attribute return type for unit
+  // attributes
+  // TODO: Method will be removed 2023-06-23/
+  auto emitAttrWithStorageTypeDeprecated =
+      [&](StringRef name, StringRef deprecatedName, StringRef attrName,
+          Attribute attr) {
+        auto *method =
+            opClass.addMethod(attr.getStorageType(), deprecatedName + "Attr");
+        std::string deprecatedMessage = std::string(deprecatedName) +
+                                        " is deprecated. Use " +
+                                        std::string(name) + " instead.";
+        method->setDeprecated(deprecatedMessage);
+        if (!method)
+          return;
+        method->body() << formatv(
+            "  return {0}.{1}<{2}>();", emitHelper.getAttr(attrName),
+            attr.isOptional() || attr.hasDefaultValue() ? "dyn_cast_or_null"
+                                                        : "cast",
+            attr.getStorageType());
+      };
+
   for (const NamedAttribute &namedAttr : op.getAttributes()) {
-    std::string name = op.getGetterName(namedAttr.name);
+    std::string name = op.getGetterNameAttr(namedAttr);
     if (namedAttr.attr.isDerivedAttr()) {
       emitDerivedAttr(name, namedAttr.attr);
+
+      // Emit deprecated derived attribute for unit attributes
+      // TODO: Method will be removed 2023-06-23/
+      if (namedAttr.attr.getBaseAttr().getAttrDefName() == "UnitAttr") {
+        std::string deprecatedName = op.getGetterName(namedAttr.name);
+        emitDerivedAttrDeprecated(name, deprecatedName, namedAttr.attr);
+      }
     } else {
       emitAttrWithStorageType(name, namedAttr.name, namedAttr.attr);
       emitAttrGetterWithReturnType(fctx, opClass, op, name, namedAttr.attr);
+
+      // Emit deprecated attribute for unit attributes
+      // TODO: Method will be removed 2023-06-23/
+      if (namedAttr.attr.getBaseAttr().getAttrDefName() == "UnitAttr") {
+        std::string deprecatedName = op.getGetterName(namedAttr.name);
+        emitAttrWithStorageTypeDeprecated(name, deprecatedName, namedAttr.name,
+                                          namedAttr.attr);
+        emitAttrGetterWithReturnTypeDeprecated(fctx, opClass, op, name,
+                                               deprecatedName, namedAttr.attr);
+      }
     }
   }
 
   auto derivedAttrs = make_filter_range(op.getAttributes(),
                                         [](const NamedAttribute &namedAttr) {
                                           return namedAttr.attr.isDerivedAttr();
                                         });
   if (derivedAttrs.empty())
// Generate method to materialize derived attributes as a DictionaryAttr.
     auto nonMaterializable =
         make_filter_range(derivedAttrs, [](const NamedAttribute &namedAttr) {
           return namedAttr.attr.getConvertFromStorageCall().empty();
         });
     if (!nonMaterializable.empty()) {
       std::string attrs;
       llvm::raw_string_ostream os(attrs);
       interleaveComma(nonMaterializable, os, [&](const NamedAttribute &attr) {
-        os << op.getGetterName(attr.name);
+        os << op.getGetterNameAttr(attr);
       });
       PrintWarning(
           op.getLoc(),
           formatv(
               "op has non-materializable derived attributes '{0}', skipping",
               os.str()));
       body << formatv("  emitOpError(\"op has non-materializable derived "
                       "attributes '{0}'\");\n",
for (const NamedAttribute &namedAttr : op.getAttributes()) {
     emitAttrWithStorageType(setterName, getterName, namedAttr.attr);
     emitAttrWithReturnType(setterName, getterName, namedAttr.attr);
   }
 }
 
 void OpEmitter::genOptionalAttrRemovers() {
   // Generate methods for removing optional attributes, instead of having to
   // use the string interface. Enables better compile time verification.
-  auto emitRemoveAttr = [&](StringRef name, bool useProperties) {
-    auto upperInitial = name.take_front().upper();
-    auto *method = opClass.addMethod("::mlir::Attribute",
-                                     op.getRemoverName(name) + "Attr");
+  auto emitRemoveAttr = [&](const NamedAttribute &namedAttr,
+                            bool useProperties) {
+    auto upperInitial = namedAttr.name.take_front().upper();
+    auto *method = opClass.addMethod(
+        "::mlir::Attribute", op.getRemoverName(namedAttr.name) + "Attr");
     if (!method)
       return;
     if (useProperties) {
       method->body() << formatv(R"(
     auto &attr = getProperties().{0};
     attr = {{};
     return attr;
 )",
-                                name);
+                                namedAttr.name);
       return;
     }
     method->body() << formatv("return (*this)->removeAttr({0}AttrName());",
-                              op.getGetterName(name));
+                              op.getGetterName(namedAttr.name));
   };
 
   for (const NamedAttribute &namedAttr : op.getAttributes())
     if (namedAttr.attr.isOptional())
-      emitRemoveAttr(namedAttr.name,
-                     op.getDialect().usePropertiesForAttributes());
+      emitRemoveAttr(namedAttr, op.getDialect().usePropertiesForAttributes());
 }
 
 // Generates the code to compute the start and end index of an operand or result
 // range.
 template <typename RangeT>
 static void generateValueRangeStartAndEnd(
     Class &opClass, bool isGenericAdaptorBase, StringRef methodName,
     int numVariadic, int numNonVariadic, StringRef rangeSizeCall,
if (!m)
     return;
 
   auto &body = m->body();
 
   // Push all result types to the operation state
   std::string resultType;
   const auto &namedAttr = op.getAttribute(0);
 
-  body << "  auto attrName = " << op.getGetterName(namedAttr.name)
-       << "AttrName(" << builderOpState
+  body << "  auto attrName = " << op.getGetterName(namedAttr.name) << "AttrName("
+       << builderOpState
        << ".name);\n"
           "  for (auto attr : attributes) {\n"
           "    if (attr.getName() != attrName) continue;\n";
   if (namedAttr.attr.isTypeAttr()) {
     resultType = "attr.getValue().cast<::mlir::TypeAttr>().getValue()";
   } else {
     resultType = "attr.getValue().cast<::mlir::TypedAttr>().getType()";
   }
if (useProperties) {
     ERROR_IF_PRUNED(m, "Adaptor::getAttributes", op);
     m->body() << "  return odsAttrs;";
   }
   for (auto &namedAttr : op.getAttributes()) {
     const auto &name = namedAttr.name;
     const auto &attr = namedAttr.attr;
     if (attr.isDerivedAttr())
       continue;
-    std::string emitName = op.getGetterName(name);
+    std::string emitName = op.getGetterNameAttr(namedAttr);
     emitAttrWithStorageType(name, emitName, attr);
     emitAttrGetterWithReturnType(fctx, genericAdaptorBase, op, emitName, attr);
   }
 
   unsigned numRegions = op.getNumRegions();
   for (unsigned i = 0; i < numRegions; ++i) {
     const auto &region = op.getRegion(i);
     if (region.name.empty())

mlir/tools/mlir-tblgen/OpFormatGen.cpp

for (const NamedAttribute &namedAttr : op.getAttributes()) {
     if (!attr.isDerivedAttr() && attr.hasDefaultValue()) {
       const StringRef &name = namedAttr.name;
       FmtContext fctx;
       fctx.withBuilder("odsBuilder");
       std::string defaultValue = std::string(
           tgfmt(attr.getConstBuilderTemplate(), &fctx, attr.getDefaultValue()));
       body << "  {\n";
       body << "     ::mlir::Builder odsBuilder(getContext());\n";
-      body << "     ::mlir::Attribute attr = " << op.getGetterName(name)
-           << "Attr();\n";
+      body << "     ::mlir::Attribute attr = "
+           << op.getGetterNameAttr(namedAttr) << "Attr();\n";
       body << "     if(attr && (attr == " << defaultValue << "))\n";
       body << "       elidedAttrs.push_back(\"" << name << "\");\n";
       body << "  }\n";
     }
   }
   body << "  _odsPrinter.printOptionalAttrDict"
        << (withKeyword ? "WithKeyword" : "")
        << "((*this)->getAttrs(), elidedAttrs);\n";
static void genSpacePrinter(bool value, MethodBody &body, bool &shouldEmitSpace,
   shouldEmitSpace = false;
 }
 
 /// Generate the printer for a custom directive parameter.
 static void genCustomDirectiveParameterPrinter(FormatElement *element,
                                                const Operator &op,
                                                MethodBody &body) {
   if (auto *attr = dyn_cast<AttributeVariable>(element)) {
-    body << op.getGetterName(attr->getVar()->name) << "Attr()";
+    body << op.getGetterNameAttr(*attr->getVar()) << "Attr()";
 
   } else if (isa<AttrDictDirective>(element)) {
     body << "getOperation()->getAttrDictionary()";
 
   } else if (auto *operand = dyn_cast<OperandVariable>(element)) {
     body << op.getGetterName(operand->getVar()->name) << "()";
 
   } else if (auto *region = dyn_cast<RegionVariable>(element)) {

 static void genEnumAttrPrinter(const NamedAttribute *var, const Operator &op,
                                MethodBody &body) {
   Attribute baseAttr = var->attr.getBaseAttr();
   const EnumAttr &enumAttr = cast<EnumAttr>(baseAttr);
   std::vector<EnumAttrCase> cases = enumAttr.getAllCases();
 
   body << llvm::formatv(enumAttrBeginPrinterCode,
                         (var->attr.isOptional() ? "*" : "") +
-                            op.getGetterName(var->name),
+                            op.getGetterNameAttr(*var),
                         enumAttr.getSymbolToStringFnName());
 
   // Get a string containing all of the cases that can't be represented with a
   // keyword.
   BitVector nonKeywordCases(cases.size());
   for (auto it : llvm::enumerate(cases)) {
     if (!canFormatStringAsKeyword(it.value().getStr()))
       nonKeywordCases.set(it.index());
TypeSwitch<FormatElement *>(anchor)
       .Case<TypeDirective>([&](TypeDirective *element) {
         genOptionalGroupPrinterAnchor(element->getArg(), op, body);
       })
       .Case<FunctionalTypeDirective>([&](FunctionalTypeDirective *element) {
         genOptionalGroupPrinterAnchor(element->getInputs(), op, body);
       })
       .Case<AttributeVariable>([&](AttributeVariable *element) {
         Attribute attr = element->getVar()->attr;
-        body << op.getGetterName(element->getVar()->name) << "Attr()";
+        body << op.getGetterNameAttr(*element->getVar()) << "Attr()";
         if (attr.isOptional())
           return; // done
         if (attr.hasDefaultValue()) {
           // Consider a default-valued attribute as present if it's not the
           // default value.
           FmtContext fctx;
           fctx.withBuilder("::mlir::OpBuilder((*this)->getContext())");
-          body << " && " << op.getGetterName(element->getVar()->name)
+          body << " && " << op.getGetterNameAttr(*element->getVar())
                << "Attr() != "
                << tgfmt(attr.getConstBuilderTemplate(), &fctx,
                         attr.getDefaultValue());
           return;
         }
         llvm_unreachable("attribute must be optional or default-valued");
       });
 }
for (auto clause : oilist->getClauses()) {
 
       SmallVector<VariableElement *> vars;
       for (FormatElement *el : pelement)
         collect(el, vars);
       body << "  if (false";
       for (VariableElement *var : vars) {
         TypeSwitch<FormatElement *>(var)
             .Case([&](AttributeVariable *attrEle) {
-              body << " || " << op.getGetterName(attrEle->getVar()->name)
+              body << " || " << op.getGetterNameAttr(*attrEle->getVar())
                    << "Attr()";
             })
             .Case([&](OperandVariable *ele) {
               if (ele->getVar()->isVariadic()) {
                 body << " || " << op.getGetterName(ele->getVar()->name)
                      << "().size()";
               } else {
                 body << " || " << op.getGetterName(ele->getVar()->name) << "()";
if (auto *attr = dyn_cast<AttributeVariable>(element)) {
     const NamedAttribute *var = attr->getVar();
 
     // If we are formatting as an enum, symbolize the attribute as a string.
     if (canFormatEnumAttr(var))
       return genEnumAttrPrinter(var, op, body);
 
     // If we are formatting as a symbol name, handle it as a symbol name.
     if (shouldFormatSymbolNameAttr(var)) {
-      body << "  _odsPrinter.printSymbolName(" << op.getGetterName(var->name)
+      body << "  _odsPrinter.printSymbolName(" << op.getGetterNameAttr(*var)
            << "Attr().getValue());\n";
       return;
     }
 
     // Elide the attribute type if it is buildable.
     if (attr->getTypeBuilder())
       body << "  _odsPrinter.printAttributeWithoutType("
-           << op.getGetterName(var->name) << "Attr());\n";
+           << op.getGetterNameAttr(*var) << "Attr());\n";
     else if (attr->shouldBeQualified() ||
              var->attr.getStorageType() == "::mlir::Attribute")
-      body << "  _odsPrinter.printAttribute(" << op.getGetterName(var->name)
+      body << "  _odsPrinter.printAttribute(" << op.getGetterNameAttr(*var)
            << "Attr());\n";
     else
       body << "_odsPrinter.printStrippedAttrOrType("
-           << op.getGetterName(var->name) << "Attr());\n";
+           << op.getGetterNameAttr(*var) << "Attr());\n";
   } else if (auto *operand = dyn_cast<OperandVariable>(element)) {
     if (operand->getVar()->isVariadicOfVariadic()) {
       body << "  ::llvm::interleaveComma("
            << op.getGetterName(operand->getVar()->name)
            << "(), _odsPrinter, [&](const auto &operands) { _odsPrinter << "
               "\"(\" << operands << "
               "\")\"; });\n";
 

mlir/tools/mlir-tblgen/SPIRVUtilsGen.cpp

for (const NamedAttribute &namedAttr : srcOp.getAttributes()) {
         // For BitEnumAttr, we need to iterate over each bit to query its
         // availability spec.
         os << formatv("  for (unsigned i = 0; "
                       "i < std::numeric_limits<{0}>::digits; ++i) {{\n",
                       enumAttr.getUnderlyingType());
         os << formatv("    {0}::{1} tblgen_attrVal = this->{2}() & "
                       "static_cast<{0}::{1}>(1 << i);\n",
                       enumAttr.getCppNamespace(), enumAttr.getEnumClassName(),
-                      srcOp.getGetterName(namedAttr.name));
+                      srcOp.getGetterNameAttr(namedAttr));
         os << formatv(
             "    if (static_cast<{0}>(tblgen_attrVal) == 0) continue;\n",
             enumAttr.getUnderlyingType());
       } else {
         // For IntEnumAttr, we just need to query the value as a whole.
         os << "  {\n";
         os << formatv("    auto tblgen_attrVal = this->{0}();\n",
-                      srcOp.getGetterName(namedAttr.name));
+                      srcOp.getGetterNameAttr(namedAttr));
       }
       os << formatv("    auto tblgen_instance = {0}::{1}(tblgen_attrVal);\n",
                     enumAttr.getCppNamespace(), avail.getQueryFnName());
       os << "    if (tblgen_instance) "
          // TODO` here once ODS supports
          // dialect-specific contents so that we can use not implementing the
          // availability interface as indication of no requirements.
          << std::string(tgfmt(caseSpecs.front().second.getMergeActionCode(),