diff --git a/mlir/include/mlir/Dialect/Async/IR/AsyncOps.td b/mlir/include/mlir/Dialect/Async/IR/AsyncOps.td --- a/mlir/include/mlir/Dialect/Async/IR/AsyncOps.td +++ b/mlir/include/mlir/Dialect/Async/IR/AsyncOps.td @@ -434,7 +434,7 @@ def Async_RuntimeStoreOp : Async_Op<"runtime.store", [TypesMatchWith<"type of 'value' matches element type of 'storage'", "storage", "value", - "$_self.cast().getValueType()">]> { + "$_self.cast().valueType()">]> { let summary = "stores the value into the runtime async.value"; let description = [{ The `async.runtime.store` operation stores the value into the runtime @@ -449,7 +449,7 @@ def Async_RuntimeLoadOp : Async_Op<"runtime.load", [TypesMatchWith<"type of 'value' matches element type of 'storage'", "storage", "result", - "$_self.cast().getValueType()">]> { + "$_self.cast().valueType()">]> { let summary = "loads the value from the runtime async.value"; let description = [{ The `async.runtime.load` operation loads the value from the runtime diff --git a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td --- a/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td +++ b/mlir/include/mlir/Dialect/LLVMIR/NVVMOps.td @@ -525,9 +525,9 @@ ); } -// Returns a list of operation suffixes corresponding to possible b1 -// multiply-and-accumulate operations for all fragments which have a -// b1 type. For all other fragments, the list returned holds a list +// Returns a list of operation suffixes corresponding to possible b1 +// multiply-and-accumulate operations for all fragments which have a +// b1 type. For all other fragments, the list returned holds a list // containing the empty string. class NVVM_MMA_B1OPS frags> { list ret = !cond( @@ -546,7 +546,7 @@ # "_" # ALayout # "_" # BLayout # !if(Satfinite, "_satfinite", "") - # signature; + # signature; } /// Helper to create the mapping between the configuration and the mma.sync @@ -563,13 +563,13 @@ "if (layoutA == \"" # layoutA # "\" && layoutB == \"" # layoutB # "\" && " " m == " # op[0].m # " && n == " # op[0].n # " && k == " # op[0].k # " && \"" # op[0].ptx_elt_type # "\" == eltypeA && \"" - # op[1].ptx_elt_type # "\" == eltypeB && " + # op[1].ptx_elt_type # "\" == eltypeB && " # " \"" # op[2].ptx_elt_type # "\" == eltypeC && " # " \"" # op[3].ptx_elt_type # "\" == eltypeD " # " && (sat.hasValue() ? " # sat # " == static_cast(*sat) : true)" # !if(!ne(b1op, ""), " && (b1Op.hasValue() ? MMAB1Op::" # b1op # " == b1Op.getValue() : true)", "") # ")\n" # " return " # - MMA_SYNC_NAME.id # ";", + MMA_SYNC_NAME.id # ";", "") // if supported ) // b1op ) // sat @@ -577,7 +577,7 @@ ) // layoutA ); // all_mma_sync_ops list>> f1 = !foldl([[[""]]], - !foldl([[[[""]]]], cond0, acc, el, + !foldl([[[[""]]]], cond0, acc, el, !listconcat(acc, el)), acc1, el1, !listconcat(acc1, el1)); list> f2 = !foldl([[""]], f1, acc1, el1, !listconcat(acc1, el1)); @@ -875,32 +875,32 @@ let description = [{ The `nvvm.mma.sync` operation collectively performs the operation - `D = matmul(A, B) + C` using all threads in a warp. + `D = matmul(A, B) + C` using all threads in a warp. All the threads in the warp must execute the same `mma.sync` operation. For each possible multiplicand PTX data type, there are one or more possible instruction shapes given as "mMnNkK". The below table describes the posssibilities - as well as the types required for the operands. Note that the data type for - C (the accumulator) and D (the result) can vary independently when there are + as well as the types required for the operands. Note that the data type for + C (the accumulator) and D (the result) can vary independently when there are multiple possibilities in the "C/D Type" column. When an optional attribute cannot be immediately inferred from the types of the operands and the result during parsing or validation, an error will be raised. - `b1Op` is only relevant when the binary (b1) type is given to + `b1Op` is only relevant when the binary (b1) type is given to `multiplicandDataType`. It specifies how the multiply-and-acumulate is performed and is either `xor_popc` or `and_poc`. The default is `xor_popc`. - `intOverflowBehavior` is only relevant when the `multiplicandType` attribute + `intOverflowBehavior` is only relevant when the `multiplicandType` attribute is one of `u8, s8, u4, s4`, this attribute describes how overflow is handled in the accumulator. When the attribute is `satfinite`, the accumulator values are clamped in the int32 range on overflow. This is the default behavior. - Alternatively, accumulator behavior `wrapped` can also be specified, in + Alternatively, accumulator behavior `wrapped` can also be specified, in which case overflow wraps from one end of the range to the other. - `layoutA` and `layoutB` are required and should generally be set to + `layoutA` and `layoutB` are required and should generally be set to `#nvvm.mma_layout` and `#nvvm.mma_layout` respectively, but other combinations are possible for certain layouts according to the table below. @@ -929,12 +929,12 @@ Example: ```mlir - %128 = nvvm.mma.sync A[%120, %121, %122, %123] - B[%124, %125] - C[%126, %127] - {layoutA = #nvvm.mma_layout, - layoutB = #nvvm.mma_layout, - shape = {k = 16 : i32, m = 16 : i32, n = 8 : i32}} + %128 = nvvm.mma.sync A[%120, %121, %122, %123] + B[%124, %125] + C[%126, %127] + {layoutA = #nvvm.mma_layout, + layoutB = #nvvm.mma_layout, + shape = {k = 16 : i32, m = 16 : i32, n = 8 : i32}} : (vector<2xf16>, vector<2xf16>, vector<2xf16>) -> !llvm.struct<(vector<2xf16>, vector<2xf16>)> ``` @@ -942,7 +942,7 @@ let results = (outs LLVM_AnyStruct:$res); let arguments = (ins NVVM_MMAShapeAttr:$shape, - OptionalAttr:$b1Op, + OptionalAttr:$b1Op, OptionalAttr:$intOverflowBehavior, MMALayoutAttr:$layoutA, MMALayoutAttr:$layoutB, @@ -950,12 +950,12 @@ OptionalAttr:$multiplicandBPtxType, Variadic:$operandA, Variadic:$operandB, - Variadic:$operandC); + Variadic:$operandC); let extraClassDeclaration = !strconcat([{ static llvm::Intrinsic::ID getIntrinsicID( int64_t m, int64_t n, uint64_t k, - llvm::Optional b1Op, + llvm::Optional b1Op, llvm::Optional sat, mlir::NVVM::MMALayout layoutAEnum, mlir::NVVM::MMALayout layoutBEnum, mlir::NVVM::MMATypes eltypeAEnum, mlir::NVVM::MMATypes eltypeBEnum, @@ -979,7 +979,7 @@ }]); let builders = [ - OpBuilder<(ins "Type":$resultType, "ValueRange":$operandA, + OpBuilder<(ins "Type":$resultType, "ValueRange":$operandA, "ValueRange":$operandB, "ValueRange":$operandC, "ArrayRef":$shape, "Optional":$b1Op, "Optional":$intOverflow, @@ -990,12 +990,12 @@ string llvmBuilder = [{ auto operands = moduleTranslation.lookupValues(opInst.getOperands()); auto intId = mlir::NVVM::MmaOp::getIntrinsicID( - $shape.getM(), $shape.getN(), $shape.getK(), + $shape.m(), $shape.n(), $shape.k(), $b1Op, $intOverflowBehavior, $layoutA, $layoutB, - $multiplicandAPtxType.getValue(), + $multiplicandAPtxType.getValue(), $multiplicandBPtxType.getValue(), - op.accumPtxType(), + op.accumPtxType(), op.resultPtxType()); $res = createIntrinsicCall( diff --git a/mlir/include/mlir/Dialect/PDL/IR/PDLTypes.td b/mlir/include/mlir/Dialect/PDL/IR/PDLTypes.td --- a/mlir/include/mlir/Dialect/PDL/IR/PDLTypes.td +++ b/mlir/include/mlir/Dialect/PDL/IR/PDLTypes.td @@ -106,7 +106,7 @@ // A range of positional values of one of the provided types. class PDL_RangeOf : ContainerType, PDL_Range.predicate, - "$_self.cast<::mlir::pdl::RangeType>().getElementType()", + "$_self.cast<::mlir::pdl::RangeType>().elementType()", "range", "::mlir::pdl::RangeType">, BuildableType<"::mlir::pdl::RangeType::get(" # positionalType.builderCall # ")">; diff --git a/mlir/include/mlir/Dialect/PDLInterp/IR/PDLInterpOps.td b/mlir/include/mlir/Dialect/PDLInterp/IR/PDLInterpOps.td --- a/mlir/include/mlir/Dialect/PDLInterp/IR/PDLInterpOps.td +++ b/mlir/include/mlir/Dialect/PDLInterp/IR/PDLInterpOps.td @@ -417,7 +417,7 @@ ```mlir // Create an instance of a `foo.op` operation. %op = pdl_interp.create_operation "foo.op"(%arg0 : !pdl.value) {"attrA" = %attr0} -> (%type : !pdl.type) - + // Create an instance of a `foo.op` operation that has inferred result types // (using the InferTypeOpInterface). %op = pdl_interp.create_operation "foo.op"(%arg0 : !pdl.value) {"attrA" = %attr0} -> @@ -559,8 +559,7 @@ let builders = [ OpBuilder<(ins "Value":$range, "unsigned":$index), [{ build($_builder, $_state, - range.getType().cast().getElementType(), - range, index); + range.getType().cast().elementType(), range, index); }]>, ]; } diff --git a/mlir/include/mlir/IR/EnumAttr.td b/mlir/include/mlir/IR/EnumAttr.td --- a/mlir/include/mlir/IR/EnumAttr.td +++ b/mlir/include/mlir/IR/EnumAttr.td @@ -386,7 +386,10 @@ let returnType = enumInfo.cppNamespace # "::" # enumInfo.className; // Convert from attribute to the underlying C++ type in op getters. - let convertFromStorage = "$_self.getValue()"; + let convertFromStorage = + !if(!or(!eq(dialect.emitAccessorPrefix, kEmitAccessorPrefix_Raw), + !eq(dialect.emitAccessorPrefix, kEmitAccessorPrefix_Both)), + "$_self.value()", "$_self.getValue()"); // The enum attribute has one parameter: the C++ enum value. let parameters = (ins EnumParameter:$value); diff --git a/mlir/include/mlir/TableGen/AttrOrTypeDef.h b/mlir/include/mlir/TableGen/AttrOrTypeDef.h --- a/mlir/include/mlir/TableGen/AttrOrTypeDef.h +++ b/mlir/include/mlir/TableGen/AttrOrTypeDef.h @@ -16,6 +16,7 @@ #include "mlir/Support/LLVM.h" #include "mlir/TableGen/Builder.h" +#include "mlir/TableGen/Dialect.h" #include "mlir/TableGen/Trait.h" namespace llvm { @@ -49,8 +50,9 @@ /// AttrOrTypeDefs to parameterize them. class AttrOrTypeParameter { public: - explicit AttrOrTypeParameter(const llvm::DagInit *def, unsigned index) - : def(def), index(index) {} + explicit AttrOrTypeParameter(const llvm::DagInit *def, unsigned index, + Dialect::EmitPrefix prefixType) + : def(def), index(index), prefixType(prefixType) {} /// Returns true if the parameter is anonymous (has no name). bool isAnonymous() const; @@ -58,6 +60,9 @@ /// Get the parameter name. StringRef getName() const; + /// Get the parameter accessor names based on the dialect setting. + SmallVector getAccessorNames() const; + /// If specified, get the custom allocator code for this parameter. Optional getAllocator() const; @@ -112,6 +117,8 @@ const llvm::DagInit *def; /// The index of the parameter within the parameter list (`def`). unsigned index; + /// The accessor prefix type. + Dialect::EmitPrefix prefixType; }; //===----------------------------------------------------------------------===// diff --git a/mlir/lib/CAPI/Dialect/PDL.cpp b/mlir/lib/CAPI/Dialect/PDL.cpp --- a/mlir/lib/CAPI/Dialect/PDL.cpp +++ b/mlir/lib/CAPI/Dialect/PDL.cpp @@ -61,7 +61,7 @@ } MlirType mlirPDLRangeTypeGetElementType(MlirType type) { - return wrap(unwrap(type).cast().getElementType()); + return wrap(unwrap(type).cast().elementType()); } //===---------------------------------------------------------------------===// diff --git a/mlir/lib/CAPI/Dialect/SparseTensor.cpp b/mlir/lib/CAPI/Dialect/SparseTensor.cpp --- a/mlir/lib/CAPI/Dialect/SparseTensor.cpp +++ b/mlir/lib/CAPI/Dialect/SparseTensor.cpp @@ -49,23 +49,23 @@ } MlirAffineMap mlirSparseTensorEncodingAttrGetDimOrdering(MlirAttribute attr) { - return wrap(unwrap(attr).cast().getDimOrdering()); + return wrap(unwrap(attr).cast().dimOrdering()); } intptr_t mlirSparseTensorEncodingGetNumDimLevelTypes(MlirAttribute attr) { - return unwrap(attr).cast().getDimLevelType().size(); + return unwrap(attr).cast().dimLevelType().size(); } MlirSparseTensorDimLevelType mlirSparseTensorEncodingAttrGetDimLevelType(MlirAttribute attr, intptr_t pos) { return static_cast( - unwrap(attr).cast().getDimLevelType()[pos]); + unwrap(attr).cast().dimLevelType()[pos]); } int mlirSparseTensorEncodingAttrGetPointerBitWidth(MlirAttribute attr) { - return unwrap(attr).cast().getPointerBitWidth(); + return unwrap(attr).cast().pointerBitWidth(); } int mlirSparseTensorEncodingAttrGetIndexBitWidth(MlirAttribute attr) { - return unwrap(attr).cast().getIndexBitWidth(); + return unwrap(attr).cast().indexBitWidth(); } diff --git a/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp b/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp --- a/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp +++ b/mlir/lib/Conversion/AsyncToLLVM/AsyncToLLVM.cpp @@ -561,7 +561,7 @@ auto loc = op->getLoc(); auto i64 = rewriter.getI64Type(); - auto storedType = converter->convertType(valueType.getValueType()); + auto storedType = converter->convertType(valueType.valueType()); auto storagePtrType = LLVM::LLVMPointerType::get(storedType); // %Size = getelementptr %T* null, int 1 @@ -1119,7 +1119,7 @@ ConversionTarget &target) { typeConverter.addConversion([&](TokenType type) { return type; }); typeConverter.addConversion([&](ValueType type) { - Type converted = typeConverter.convertType(type.getValueType()); + Type converted = typeConverter.convertType(type.valueType()); return converted ? ValueType::get(converted) : converted; }); diff --git a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp --- a/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp +++ b/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp @@ -435,7 +435,7 @@ return parallelOp.emitOpError() << "expected mapping attribute for lowering to GPU"; Value newIndex; - gpu::Processor processor = annotation.getProcessor(); + gpu::Processor processor = annotation.processor(); if (isMappedToProcessor(processor)) { // Use the corresponding thread/grid index as replacement for the loop iv. @@ -450,11 +450,11 @@ rewriter.getAffineDimExpr(0) * rewriter.getAffineSymbolExpr(0) + rewriter.getAffineSymbolExpr(1)); newIndex = rewriter.create( - loc, annotation.getMap().compose(lowerAndStep), + loc, annotation.map().compose(lowerAndStep), ValueRange{operand, step, lowerBound}); // If there was also a bound, insert that, too. // TODO: Check that we do not assign bounds twice. - if (annotation.getBound()) { + if (annotation.bound()) { // We pass as the single operand to the bound-map the number of // iterations, which is (upperBound - lowerBound) ceilDiv step. To // support inner loops with dynamic upper bounds (as generated by e.g. @@ -494,7 +494,7 @@ ((rewriter.getAffineDimExpr(0) - rewriter.getAffineSymbolExpr(0)) .ceilDiv(rewriter.getAffineSymbolExpr(1)))); Value launchBound = rewriter.create( - loc, annotation.getBound().compose(stepMap), + loc, annotation.bound().compose(stepMap), ValueRange{ ensureLaunchIndependent( cloningMap.lookupOrDefault(upperBound)), diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalg.cpp @@ -147,8 +147,8 @@ cast(op).quantization_info()) { auto quantizationInfo = cast(op).quantization_info(); int32_t inputBitWidth = elementTy.getIntOrFloatBitWidth(); - int64_t inZp = quantizationInfo.getValue().getInput_zp(); - int64_t outZp = quantizationInfo.getValue().getOutput_zp(); + int64_t inZp = quantizationInfo.getValue().input_zp(); + int64_t outZp = quantizationInfo.getValue().output_zp(); // Compute the maximum value that can occur in the intermediate buffer. int64_t zpAdd = inZp + outZp; @@ -1847,7 +1847,7 @@ } else if (elementTy.isa() && !padOp.quantization_info()) { constantAttr = rewriter.getIntegerAttr(elementTy, 0); } else if (elementTy.isa() && padOp.quantization_info()) { - int64_t value = padOp.quantization_info().getValue().getInput_zp(); + int64_t value = padOp.quantization_info().getValue().input_zp(); constantAttr = rewriter.getIntegerAttr(elementTy, value); } if (constantAttr) diff --git a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp --- a/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp +++ b/mlir/lib/Conversion/TosaToLinalg/TosaToLinalgNamed.cpp @@ -202,7 +202,7 @@ if (isQuantized) { auto quantizationInfo = op->getAttr("quantization_info").cast(); - int64_t iZp = quantizationInfo.getInput_zp(); + int64_t iZp = quantizationInfo.input_zp(); int64_t intMin = APInt::getSignedMinValue(inputETy.getIntOrFloatBitWidth()) @@ -274,8 +274,8 @@ if (isQuantized) { auto quantizationInfo = op->getAttr("quantization_info").cast(); - auto iZp = rewriter.getI32IntegerAttr(quantizationInfo.getInput_zp()); - auto kZp = rewriter.getI32IntegerAttr(quantizationInfo.getWeight_zp()); + auto iZp = rewriter.getI32IntegerAttr(quantizationInfo.input_zp()); + auto kZp = rewriter.getI32IntegerAttr(quantizationInfo.weight_zp()); auto iZpVal = rewriter.create(loc, iZp); auto kZpVal = rewriter.create(loc, kZp); @@ -366,8 +366,8 @@ if (isQuantized) { auto quantizationInfo = op->getAttr("quantization_info").cast(); - iZp = rewriter.getI32IntegerAttr(quantizationInfo.getInput_zp()); - kZp = rewriter.getI32IntegerAttr(quantizationInfo.getWeight_zp()); + iZp = rewriter.getI32IntegerAttr(quantizationInfo.input_zp()); + kZp = rewriter.getI32IntegerAttr(quantizationInfo.weight_zp()); } auto weightShape = weightTy.getShape(); @@ -378,7 +378,7 @@ if (isQuantized) { auto quantizationInfo = op->getAttr("quantization_info").cast(); - int64_t iZp = quantizationInfo.getInput_zp(); + int64_t iZp = quantizationInfo.input_zp(); int64_t intMin = APInt::getSignedMinValue(inputETy.getIntOrFloatBitWidth()) @@ -542,9 +542,9 @@ auto quantizationInfo = op.quantization_info().getValue(); auto aZp = rewriter.create( - loc, rewriter.getI32IntegerAttr(quantizationInfo.getA_zp())); + loc, rewriter.getI32IntegerAttr(quantizationInfo.a_zp())); auto bZp = rewriter.create( - loc, rewriter.getI32IntegerAttr(quantizationInfo.getB_zp())); + loc, rewriter.getI32IntegerAttr(quantizationInfo.b_zp())); rewriter.replaceOpWithNewOp( op, TypeRange{op.getType()}, ValueRange{adaptor.a(), adaptor.b(), aZp, bZp}, zeroTensor); @@ -652,9 +652,9 @@ auto quantizationInfo = op.quantization_info().getValue(); auto inputZp = rewriter.create( - loc, rewriter.getI32IntegerAttr(quantizationInfo.getInput_zp())); + loc, rewriter.getI32IntegerAttr(quantizationInfo.input_zp())); auto outputZp = rewriter.create( - loc, rewriter.getI32IntegerAttr(quantizationInfo.getWeight_zp())); + loc, rewriter.getI32IntegerAttr(quantizationInfo.weight_zp())); Value matmul = rewriter .create( @@ -892,8 +892,7 @@ if (op.quantization_info()) { auto quantizationInfo = op.quantization_info().getValue(); auto inputZp = rewriter.create( - loc, - b.getIntegerAttr(accETy, quantizationInfo.getInput_zp())); + loc, b.getIntegerAttr(accETy, quantizationInfo.input_zp())); Value offset = rewriter.create(loc, accETy, countI, inputZp); poolVal = @@ -930,7 +929,7 @@ auto quantizationInfo = op.quantization_info().getValue(); auto outputZp = rewriter.create( loc, b.getIntegerAttr(scaled.getType(), - quantizationInfo.getOutput_zp())); + quantizationInfo.output_zp())); scaled = rewriter.create(loc, scaled, outputZp) .getResult(); } diff --git a/mlir/lib/Dialect/Async/IR/Async.cpp b/mlir/lib/Dialect/Async/IR/Async.cpp --- a/mlir/lib/Dialect/Async/IR/Async.cpp +++ b/mlir/lib/Dialect/Async/IR/Async.cpp @@ -38,7 +38,7 @@ // parent `async.execute` operation. auto executeOp = (*this)->getParentOfType(); auto types = llvm::map_range(executeOp.results(), [](const OpResult &result) { - return result.getType().cast().getValueType(); + return result.getType().cast().valueType(); }); if (getOperandTypes() != types) @@ -67,7 +67,7 @@ bool ExecuteOp::areTypesCompatible(Type lhs, Type rhs) { const auto getValueOrTokenType = [](Type type) { if (auto value = type.dyn_cast()) - return value.getValueType(); + return value.valueType(); return type; }; return getValueOrTokenType(lhs) == getValueOrTokenType(rhs); @@ -114,8 +114,7 @@ Block &bodyBlock = bodyRegion->front(); for (Value operand : operands) { auto valueType = operand.getType().dyn_cast(); - bodyBlock.addArgument(valueType ? valueType.getValueType() - : operand.getType(), + bodyBlock.addArgument(valueType ? valueType.valueType() : operand.getType(), operand.getLoc()); } @@ -190,7 +189,7 @@ return failure(); auto valueTy = valueTypes.back().dyn_cast(); - unwrappedArgs.back().type = valueTy ? valueTy.getValueType() : Type(); + unwrappedArgs.back().type = valueTy ? valueTy.valueType() : Type(); return success(); }; @@ -229,7 +228,7 @@ LogicalResult ExecuteOp::verifyRegions() { // Unwrap async.execute value operands types. auto unwrappedTypes = llvm::map_range(operands(), [](Value operand) { - return operand.getType().cast().getValueType(); + return operand.getType().cast().valueType(); }); // Verify that unwrapped argument types matches the body region arguments. @@ -281,7 +280,7 @@ // Add unwrapped async.value type to the returned values types. if (auto valueType = operand.getType().dyn_cast()) - result.addTypes(valueType.getValueType()); + result.addTypes(valueType.valueType()); } static ParseResult parseAwaitResultType(OpAsmParser &parser, Type &operandType, @@ -291,7 +290,7 @@ // Add unwrapped async.value type to the returned values types. if (auto valueType = operandType.dyn_cast()) - resultType = valueType.getValueType(); + resultType = valueType.valueType(); return success(); } @@ -310,10 +309,10 @@ // Awaiting on a value unwraps the async value type. if (auto value = argType.dyn_cast()) { - if (*getResultType() != value.getValueType()) + if (*getResultType() != value.valueType()) return emitOpError() << "result type " << *getResultType() << " does not match async value type " - << value.getValueType(); + << value.valueType(); } return success(); @@ -335,7 +334,7 @@ void ValueType::print(AsmPrinter &printer) const { printer << "<"; - printer.printType(getValueType()); + printer.printType(valueType()); printer << '>'; } diff --git a/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp b/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp --- a/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp +++ b/mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp @@ -503,7 +503,7 @@ getReplacementValue(AwaitOp op, Value operand, ConversionPatternRewriter &rewriter) const override { // Load from the async value storage. - auto valueType = operand.getType().cast().getValueType(); + auto valueType = operand.getType().cast().valueType(); return rewriter.create(op->getLoc(), valueType, operand); } }; diff --git a/mlir/lib/Dialect/EmitC/IR/EmitC.cpp b/mlir/lib/Dialect/EmitC/IR/EmitC.cpp --- a/mlir/lib/Dialect/EmitC/IR/EmitC.cpp +++ b/mlir/lib/Dialect/EmitC/IR/EmitC.cpp @@ -209,7 +209,7 @@ void emitc::OpaqueAttr::print(AsmPrinter &printer) const { printer << "<\""; - llvm::printEscapedString(getValue(), printer.getStream()); + llvm::printEscapedString(value(), printer.getStream()); printer << "\">"; } @@ -245,6 +245,6 @@ void emitc::OpaqueType::print(AsmPrinter &printer) const { printer << "<\""; - llvm::printEscapedString(getValue(), printer.getStream()); + llvm::printEscapedString(value(), printer.getStream()); printer << "\">"; } diff --git a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp --- a/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/AsyncRegionRewriter.cpp @@ -150,7 +150,7 @@ [](Type type) { // Extract value type from !async.value. if (auto valueType = type.dyn_cast()) - return valueType.getValueType(); + return valueType.valueType(); assert(type.isa() && "expected token type"); return type; }); @@ -297,8 +297,7 @@ if (result.use_empty() || result.hasOneUse()) return false; auto valueType = result.getType().dyn_cast(); - return valueType && - valueType.getValueType().isa(); + return valueType && valueType.valueType().isa(); }); if (multiUseResults.empty()) return; diff --git a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp --- a/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp +++ b/mlir/lib/Dialect/GPU/Transforms/ParallelLoopMapper.cpp @@ -31,7 +31,7 @@ // Verify that each processor is mapped to only once. llvm::DenseSet specifiedMappings; for (auto dimAttr : mapping) { - gpu::Processor processor = dimAttr.getProcessor(); + gpu::Processor processor = dimAttr.processor(); if (processor != gpu::Processor::Sequential && specifiedMappings.count(processor)) return ploopOp.emitError( diff --git a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp --- a/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp +++ b/mlir/lib/Dialect/LLVMIR/IR/NVVMDialect.cpp @@ -355,8 +355,8 @@ auto s32x2StructTy = LLVM::LLVMStructType::getLiteral(context, {i32Ty, i32Ty}); - std::array mmaShape{shapeAttr().getM(), shapeAttr().getN(), - shapeAttr().getK()}; + std::array mmaShape{shapeAttr().m(), shapeAttr().n(), + shapeAttr().k()}; // These variables define the set of allowed data types for matrices A, B, C, // and result. diff --git a/mlir/lib/Dialect/Linalg/Transforms/SparseTensorRewriting.cpp b/mlir/lib/Dialect/Linalg/Transforms/SparseTensorRewriting.cpp --- a/mlir/lib/Dialect/Linalg/Transforms/SparseTensorRewriting.cpp +++ b/mlir/lib/Dialect/Linalg/Transforms/SparseTensorRewriting.cpp @@ -37,7 +37,7 @@ static bool isSparseTensor(OpOperand *op) { if (auto enc = getSparseTensorEncoding(op->get().getType())) { ArrayRef dimTypes = - enc.getDimLevelType(); + enc.dimLevelType(); for (auto dimType : dimTypes) if (dimType == SparseTensorEncodingAttr::DimLevelType::Compressed) return true; // at least one compressed diff --git a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp --- a/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp +++ b/mlir/lib/Dialect/OpenMP/IR/OpenMPDialect.cpp @@ -113,7 +113,7 @@ template static ParseResult parseClauseAttr(AsmParser &parser, ClauseAttr &attr) { - using ClauseT = decltype(std::declval().getValue()); + using ClauseT = decltype(std::declval().value()); StringRef enumStr; SMLoc loc = parser.getCurrentLocation(); if (parser.parseKeyword(&enumStr)) @@ -127,7 +127,7 @@ template void printClauseAttr(OpAsmPrinter &p, Operation *op, ClauseAttr attr) { - p << stringifyEnum(attr.getValue()); + p << stringifyEnum(attr.value()); } //===----------------------------------------------------------------------===// @@ -282,11 +282,11 @@ ScheduleModifierAttr modifier, UnitAttr simd, Value scheduleChunkVar, Type scheduleChunkType) { - p << stringifyClauseScheduleKind(schedAttr.getValue()); + p << stringifyClauseScheduleKind(schedAttr.value()); if (scheduleChunkVar) p << " = " << scheduleChunkVar << " : " << scheduleChunkVar.getType(); if (modifier) - p << ", " << stringifyScheduleModifier(modifier.getValue()); + p << ", " << stringifyScheduleModifier(modifier.value()); if (simd) p << ", simd"; } diff --git a/mlir/lib/Dialect/PDL/IR/PDLTypes.cpp b/mlir/lib/Dialect/PDL/IR/PDLTypes.cpp --- a/mlir/lib/Dialect/PDL/IR/PDLTypes.cpp +++ b/mlir/lib/Dialect/PDL/IR/PDLTypes.cpp @@ -85,7 +85,7 @@ void RangeType::print(AsmPrinter &printer) const { printer << "<"; - (void)generatedTypePrinter(getElementType(), printer); + (void)generatedTypePrinter(elementType(), printer); printer << ">"; } diff --git a/mlir/lib/Dialect/PDLInterp/IR/PDLInterp.cpp b/mlir/lib/Dialect/PDLInterp/IR/PDLInterp.cpp --- a/mlir/lib/Dialect/PDLInterp/IR/PDLInterp.cpp +++ b/mlir/lib/Dialect/PDLInterp/IR/PDLInterp.cpp @@ -147,8 +147,7 @@ // FIXME: Allow passing in a proper location for the loop variable. auto rangeType = range.getType().cast(); state.regions.front()->emplaceBlock(); - state.regions.front()->addArgument(rangeType.getElementType(), - state.location); + state.regions.front()->addArgument(rangeType.elementType(), state.location); } } diff --git a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp --- a/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp +++ b/mlir/lib/Dialect/SPIRV/IR/TargetAndABI.cpp @@ -145,7 +145,7 @@ DenseIntElementsAttr spirv::lookupLocalWorkGroupSize(Operation *op) { if (auto entryPoint = spirv::lookupEntryPointABI(op)) - return entryPoint.getLocal_size(); + return entryPoint.local_size(); return {}; } diff --git a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp --- a/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp +++ b/mlir/lib/Dialect/SPIRV/Transforms/LowerABIAttributesPass.cpp @@ -135,7 +135,7 @@ funcOp.getLoc(), executionModel.getValue(), funcOp, interfaceVars); // Specifies the spv.ExecutionModeOp. - auto localSizeAttr = entryPointAttr.getLocal_size(); + auto localSizeAttr = entryPointAttr.local_size(); if (localSizeAttr) { auto values = localSizeAttr.getValues(); SmallVector localSize(values); diff --git a/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp b/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp --- a/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp +++ b/mlir/lib/Dialect/SparseTensor/IR/SparseTensorDialect.cpp @@ -118,8 +118,8 @@ void SparseTensorEncodingAttr::print(AsmPrinter &printer) const { // Print the struct-like storage in dictionary fashion. printer << "<{ dimLevelType = [ "; - for (unsigned i = 0, e = getDimLevelType().size(); i < e; i++) { - switch (getDimLevelType()[i]) { + for (unsigned i = 0, e = dimLevelType().size(); i < e; i++) { + switch (dimLevelType()[i]) { case DimLevelType::Dense: printer << "\"dense\""; break; @@ -134,10 +134,10 @@ printer << ", "; } printer << " ]"; - if (getDimOrdering()) - printer << ", dimOrdering = affine_map<" << getDimOrdering() << ">"; - printer << ", pointerBitWidth = " << getPointerBitWidth() - << ", indexBitWidth = " << getIndexBitWidth() << " }>"; + if (dimOrdering()) + printer << ", dimOrdering = affine_map<" << dimOrdering() << ">"; + printer << ", pointerBitWidth = " << pointerBitWidth() + << ", indexBitWidth = " << indexBitWidth() << " }>"; } LogicalResult SparseTensorEncodingAttr::verify( @@ -163,18 +163,18 @@ ArrayRef shape, Type elementType, function_ref emitError) const { // Check structural integrity. - if (failed(verify(emitError, getDimLevelType(), getDimOrdering(), - getPointerBitWidth(), getIndexBitWidth()))) + if (failed(verify(emitError, dimLevelType(), dimOrdering(), pointerBitWidth(), + indexBitWidth()))) return failure(); // Check integrity with tensor type specifics. Dimension ordering is optional, // but we always should have dimension level types for the full rank. unsigned size = shape.size(); if (size == 0) return emitError() << "expected non-scalar sparse tensor"; - if (getDimOrdering() && getDimOrdering().getNumResults() != size) + if (dimOrdering() && dimOrdering().getNumResults() != size) return emitError() << "expected an affine map of size " << size << " for dimension ordering"; - if (getDimLevelType().size() != size) + if (dimLevelType().size() != size) return emitError() << "expected an array of size " << size << " for dimension level types"; return success(); @@ -237,7 +237,7 @@ auto e = getSparseTensorEncoding(tensor().getType()); if (failed(isInBounds(dim(), tensor()))) return emitError("requested pointers dimension out of bounds"); - if (failed(isMatchingWidth(result(), e.getPointerBitWidth()))) + if (failed(isMatchingWidth(result(), e.pointerBitWidth()))) return emitError("unexpected type for pointers"); return success(); } @@ -246,7 +246,7 @@ auto e = getSparseTensorEncoding(tensor().getType()); if (failed(isInBounds(dim(), tensor()))) return emitError("requested indices dimension out of bounds"); - if (failed(isMatchingWidth(result(), e.getIndexBitWidth()))) + if (failed(isMatchingWidth(result(), e.indexBitWidth()))) return emitError("unexpected type for indices"); return success(); } diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.h b/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.h --- a/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.h +++ b/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.h @@ -168,14 +168,14 @@ /// overhead storage. inline Value constantPointerTypeEncoding(OpBuilder &builder, Location loc, const SparseTensorEncodingAttr &enc) { - return constantOverheadTypeEncoding(builder, loc, enc.getPointerBitWidth()); + return constantOverheadTypeEncoding(builder, loc, enc.pointerBitWidth()); } /// Generates a constant of the internal type-encoding for index overhead /// storage. inline Value constantIndexTypeEncoding(OpBuilder &builder, Location loc, const SparseTensorEncodingAttr &enc) { - return constantOverheadTypeEncoding(builder, loc, enc.getIndexBitWidth()); + return constantOverheadTypeEncoding(builder, loc, enc.indexBitWidth()); } /// Generates a constant of the internal type-encoding for primary storage. diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.cpp --- a/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/CodegenUtils.cpp @@ -60,12 +60,12 @@ OverheadType mlir::sparse_tensor::pointerOverheadTypeEncoding( const SparseTensorEncodingAttr &enc) { - return overheadTypeEncoding(enc.getPointerBitWidth()); + return overheadTypeEncoding(enc.pointerBitWidth()); } OverheadType mlir::sparse_tensor::indexOverheadTypeEncoding( const SparseTensorEncodingAttr &enc) { - return overheadTypeEncoding(enc.getIndexBitWidth()); + return overheadTypeEncoding(enc.indexBitWidth()); } Type mlir::sparse_tensor::getPointerOverheadType( diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp --- a/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/SparseTensorConversion.cpp @@ -95,7 +95,7 @@ SparseTensorEncodingAttr &enc, Value src, int64_t idx) { // Permute the index according to an optional dimension ordering. - if (AffineMap p = enc.getDimOrdering()) + if (AffineMap p = enc.dimOrdering()) idx = p.getPermutedPosition(idx); // Generate the call. StringRef name = "sparseDimSize"; @@ -197,7 +197,7 @@ SparseTensorEncodingAttr &enc, Action action, ValueRange szs, Value ptr = Value()) { Location loc = op->getLoc(); - ArrayRef dlt = enc.getDimLevelType(); + ArrayRef dlt = enc.dimLevelType(); unsigned sz = dlt.size(); // Sparsity annotations. SmallVector attrs; @@ -211,7 +211,7 @@ // default, or otherwise the "reverse" permutation of a given ordering, so // that indices can be mapped quickly to the right position. SmallVector rev(sz); - if (AffineMap p = enc.getDimOrdering()) { + if (AffineMap p = enc.dimOrdering()) { for (unsigned i = 0; i < sz; i++) rev[p.getDimPosition(i)] = constantIndex(builder, loc, i); } else { @@ -536,11 +536,11 @@ break; case SparseToSparseConversionStrategy::kDirect: useDirectConversion = true; - assert(canUseDirectConversion(encDst.getDimLevelType()) && + assert(canUseDirectConversion(encDst.dimLevelType()) && "Unsupported target for direct sparse-to-sparse conversion"); break; case SparseToSparseConversionStrategy::kAuto: - useDirectConversion = canUseDirectConversion(encDst.getDimLevelType()); + useDirectConversion = canUseDirectConversion(encDst.dimLevelType()); break; } if (useDirectConversion) { @@ -552,8 +552,8 @@ // method calls can share most parameters, while still providing // the correct sparsity information to either of them. auto enc = SparseTensorEncodingAttr::get( - op->getContext(), encDst.getDimLevelType(), encDst.getDimOrdering(), - encSrc.getPointerBitWidth(), encSrc.getIndexBitWidth()); + op->getContext(), encDst.dimLevelType(), encDst.dimOrdering(), + encSrc.pointerBitWidth(), encSrc.indexBitWidth()); newParams(rewriter, params, op, stp, enc, Action::kToCOO, sizes, src); Value coo = genNewCall(rewriter, op, params); params[3] = constantPointerTypeEncoding(rewriter, loc, encDst); @@ -585,7 +585,7 @@ op->getContext(), SmallVector( rank, SparseTensorEncodingAttr::DimLevelType::Dense), - AffineMap(), encSrc.getPointerBitWidth(), encSrc.getIndexBitWidth()); + AffineMap(), encSrc.pointerBitWidth(), encSrc.indexBitWidth()); SmallVector sizes; SmallVector params; sizesFromPtr(rewriter, sizes, op, encSrc, srcTensorTp, src); @@ -885,8 +885,8 @@ SmallVector params; sizesFromPtr(rewriter, sizes, op, encSrc, srcType, src); auto enc = SparseTensorEncodingAttr::get( - op->getContext(), encSrc.getDimLevelType(), AffineMap(), - encSrc.getPointerBitWidth(), encSrc.getIndexBitWidth()); + op->getContext(), encSrc.dimLevelType(), AffineMap(), + encSrc.pointerBitWidth(), encSrc.indexBitWidth()); newParams(rewriter, params, op, srcType, enc, Action::kToCOO, sizes, src); Value coo = genNewCall(rewriter, op, params); // Then output the tensor to external file with indices in the externally @@ -894,8 +894,7 @@ // not in that order yet (note that the sort can be dropped altogether if // external format does not care about the order at all, but here we assume // it does). - bool sort = - encSrc.getDimOrdering() && !encSrc.getDimOrdering().isIdentity(); + bool sort = encSrc.dimOrdering() && !encSrc.dimOrdering().isIdentity(); params.clear(); params.push_back(coo); params.push_back(adaptor.getOperands()[1]); diff --git a/mlir/lib/Dialect/SparseTensor/Transforms/Sparsification.cpp b/mlir/lib/Dialect/SparseTensor/Transforms/Sparsification.cpp --- a/mlir/lib/Dialect/SparseTensor/Transforms/Sparsification.cpp +++ b/mlir/lib/Dialect/SparseTensor/Transforms/Sparsification.cpp @@ -107,7 +107,7 @@ /// Helper method to apply dimension ordering permutation. static unsigned perm(const SparseTensorEncodingAttr &enc, unsigned d) { if (enc) { - auto order = enc.getDimOrdering(); + auto order = enc.dimOrdering(); if (order) { assert(order.isPermutation()); return order.getDimPosition(d); @@ -119,7 +119,7 @@ /// Helper method to translate dim level type to internal representation. static Dim toDim(const SparseTensorEncodingAttr &enc, unsigned d) { if (enc) { - SparseTensorEncodingAttr::DimLevelType tp = enc.getDimLevelType()[d]; + SparseTensorEncodingAttr::DimLevelType tp = enc.dimLevelType()[d]; if (tp == SparseTensorEncodingAttr::DimLevelType::Compressed) return Dim::kSparse; if (tp == SparseTensorEncodingAttr::DimLevelType::Singleton) diff --git a/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp b/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp --- a/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp +++ b/mlir/lib/Dialect/Tosa/IR/TosaOps.cpp @@ -347,7 +347,7 @@ } else if (elementTy.isa() && !op.quantization_info()) { constantAttr = rewriter.getIntegerAttr(elementTy, 0); } else if (elementTy.isa() && op.quantization_info()) { - auto value = op.quantization_info().getValue().getInput_zp(); + auto value = op.quantization_info().getValue().input_zp(); constantAttr = rewriter.getIntegerAttr(elementTy, value); } diff --git a/mlir/lib/Dialect/Tosa/Transforms/TosaDecomposeTransposeConv.cpp b/mlir/lib/Dialect/Tosa/Transforms/TosaDecomposeTransposeConv.cpp --- a/mlir/lib/Dialect/Tosa/Transforms/TosaDecomposeTransposeConv.cpp +++ b/mlir/lib/Dialect/Tosa/Transforms/TosaDecomposeTransposeConv.cpp @@ -214,7 +214,7 @@ weight = createOpAndInfer( rewriter, loc, UnrankedTensorType::get(weightETy), weight, weightPaddingVal, nullptr, - rewriter.getAttr(quantInfo.getWeight_zp())); + rewriter.getAttr(quantInfo.weight_zp())); } else { weight = createOpAndInfer(rewriter, loc, @@ -278,7 +278,7 @@ input = createOpAndInfer( rewriter, loc, UnrankedTensorType::get(inputETy), input, inputPaddingVal, nullptr, - rewriter.getAttr(quantInfo.getInput_zp())); + rewriter.getAttr(quantInfo.input_zp())); } else { input = createOpAndInfer(rewriter, loc, UnrankedTensorType::get(inputETy), diff --git a/mlir/lib/Rewrite/ByteCode.cpp b/mlir/lib/Rewrite/ByteCode.cpp --- a/mlir/lib/Rewrite/ByteCode.cpp +++ b/mlir/lib/Rewrite/ByteCode.cpp @@ -390,7 +390,7 @@ .Case( [](Type) { return PDLValue::Kind::Operation; }) .Case([](pdl::RangeType rangeTy) { - if (rangeTy.getElementType().isa()) + if (rangeTy.elementType().isa()) return PDLValue::Kind::TypeRange; return PDLValue::Kind::ValueRange; }) @@ -527,7 +527,7 @@ auto processRewriterValue = [&](Value val) { valueToMemIndex.try_emplace(val, index++); if (pdl::RangeType rangeType = val.getType().dyn_cast()) { - Type elementTy = rangeType.getElementType(); + Type elementTy = rangeType.elementType(); if (elementTy.isa()) valueToRangeIndex.try_emplace(val, typeRangeIndex++); else if (elementTy.isa()) @@ -600,7 +600,7 @@ // Check to see if this value is a range type. if (auto rangeTy = value.getType().dyn_cast()) { - Type eleType = rangeTy.getElementType(); + Type eleType = rangeTy.elementType(); if (eleType.isa()) defRangeIt->second.opRangeIndex = 0; else if (eleType.isa()) diff --git a/mlir/lib/TableGen/AttrOrTypeDef.cpp b/mlir/lib/TableGen/AttrOrTypeDef.cpp --- a/mlir/lib/TableGen/AttrOrTypeDef.cpp +++ b/mlir/lib/TableGen/AttrOrTypeDef.cpp @@ -59,8 +59,9 @@ // Populate the parameters. if (auto *parametersDag = def->getValueAsDag("parameters")) { + Dialect::EmitPrefix prefixType = getDialect().getEmitAccessorPrefix(); for (unsigned i = 0, e = parametersDag->getNumArgs(); i < e; ++i) - parameters.push_back(AttrOrTypeParameter(parametersDag, i)); + parameters.push_back(AttrOrTypeParameter(parametersDag, i, prefixType)); } // Verify the use of the mnemonic field. @@ -215,6 +216,19 @@ return def->getArgName(index)->getValue(); } +SmallVector AttrOrTypeParameter::getAccessorNames() const { + SmallVector names; + if (prefixType == Dialect::EmitPrefix::Both || + prefixType == Dialect::EmitPrefix::Raw) + names.push_back(getName().str()); + if (prefixType == Dialect::EmitPrefix::Both || + prefixType == Dialect::EmitPrefix::Prefixed) { + names.push_back("get" + llvm::convertToCamelFromSnakeCase( + getName(), /*capitalizeFirst=*/true)); + } + return names; +} + Optional AttrOrTypeParameter::getAllocator() const { return getDefValue("allocator"); } diff --git a/mlir/lib/Target/Cpp/TranslateToCpp.cpp b/mlir/lib/Target/Cpp/TranslateToCpp.cpp --- a/mlir/lib/Target/Cpp/TranslateToCpp.cpp +++ b/mlir/lib/Target/Cpp/TranslateToCpp.cpp @@ -191,7 +191,7 @@ if (emitter.shouldDeclareVariablesAtTop()) { // Skip the assignment if the emitc.constant has no value. if (auto oAttr = value.dyn_cast()) { - if (oAttr.getValue().empty()) + if (oAttr.value().empty()) return success(); } @@ -202,7 +202,7 @@ // Emit a variable declaration for an emitc.constant op without value. if (auto oAttr = value.dyn_cast()) { - if (oAttr.getValue().empty()) + if (oAttr.value().empty()) // The semicolon gets printed by the emitOperation function. return emitter.emitVariableDeclaration(result, /*trailingSemicolon=*/false); @@ -806,7 +806,7 @@ // Print opaque attributes. if (auto oAttr = attr.dyn_cast()) { - os << oAttr.getValue(); + os << oAttr.value(); return success(); } @@ -1003,11 +1003,11 @@ if (auto tType = type.dyn_cast()) return emitTupleType(loc, tType.getTypes()); if (auto oType = type.dyn_cast()) { - os << oType.getValue(); + os << oType.value(); return success(); } if (auto pType = type.dyn_cast()) { - if (failed(emitType(loc, pType.getPointee()))) + if (failed(emitType(loc, pType.pointee()))) return failure(); os << "*"; return success(); diff --git a/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml b/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml --- a/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml +++ b/mlir/test/mlir-linalg-ods-gen/test-linalg-ods-yaml-gen.yaml @@ -92,7 +92,7 @@ # IMPL-NEXT: return attr.getName() == "cast"; }); # IMPL-NEXT: if (castIter != attrs.end()) { # IMPL-NEXT: if (auto attr = castIter->getValue().dyn_cast()) -# IMPL-NEXT: castVal = attr.getValue(); +# IMPL-NEXT: castVal = attr.value(); # IMPL-NEXT: } # IMPL: Value [[VAL0:[a-z0-9]+]] = helper.constant("42 : i64"); diff --git a/mlir/test/mlir-tblgen/attr-or-type-format.td b/mlir/test/mlir-tblgen/attr-or-type-format.td --- a/mlir/test/mlir-tblgen/attr-or-type-format.td +++ b/mlir/test/mlir-tblgen/attr-or-type-format.td @@ -9,6 +9,7 @@ let name = "TestDialect"; let cppNamespace = "::test"; let useDefaultTypePrinterParser = 0; + let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed; } class TestAttr : AttrDef; diff --git a/mlir/test/mlir-tblgen/attrdefs.td b/mlir/test/mlir-tblgen/attrdefs.td --- a/mlir/test/mlir-tblgen/attrdefs.td +++ b/mlir/test/mlir-tblgen/attrdefs.td @@ -33,10 +33,11 @@ // DEF: return {}; def Test_Dialect: Dialect { -// DECL-NOT: TestDialect -// DEF-NOT: TestDialect - let name = "TestDialect"; - let cppNamespace = "::test"; + // DECL-NOT: TestDialect + // DEF-NOT: TestDialect + let name = "TestDialect"; + let cppNamespace = "::test"; + let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed; } class TestAttr : AttrDef { } diff --git a/mlir/test/mlir-tblgen/typedefs.td b/mlir/test/mlir-tblgen/typedefs.td --- a/mlir/test/mlir-tblgen/typedefs.td +++ b/mlir/test/mlir-tblgen/typedefs.td @@ -34,9 +34,10 @@ // DEF: return {}; def Test_Dialect: Dialect { -// DECL-NOT: TestDialect + // DECL-NOT: TestDialect let name = "TestDialect"; let cppNamespace = "::test"; + let emitAccessorPrefix = kEmitAccessorPrefix_Prefixed; } class TestType : TypeDef { } diff --git a/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp b/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp --- a/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp +++ b/mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp @@ -1027,7 +1027,7 @@ return attr.getName() == "{1}"; }); if ({1}Iter != attrs.end()) {{ if (auto attr = {1}Iter->getValue().dyn_cast<{0}Attr>()) - {1}Val = attr.getValue(); + {1}Val = attr.value(); } )FMT"; std::string enumName = convertOperandKindToEnumName(arg.kind); diff --git a/mlir/tools/mlir-tblgen/AttrOrTypeDefGen.cpp b/mlir/tools/mlir-tblgen/AttrOrTypeDefGen.cpp --- a/mlir/tools/mlir-tblgen/AttrOrTypeDefGen.cpp +++ b/mlir/tools/mlir-tblgen/AttrOrTypeDefGen.cpp @@ -7,16 +7,12 @@ //===----------------------------------------------------------------------===// #include "AttrOrTypeFormatGen.h" -#include "mlir/Support/LogicalResult.h" #include "mlir/TableGen/AttrOrTypeDef.h" #include "mlir/TableGen/Class.h" #include "mlir/TableGen/CodeGenHelpers.h" #include "mlir/TableGen/Format.h" #include "mlir/TableGen/GenInfo.h" #include "mlir/TableGen/Interfaces.h" -#include "llvm/ADT/Sequence.h" -#include "llvm/ADT/SetVector.h" -#include "llvm/ADT/SmallSet.h" #include "llvm/ADT/StringSet.h" #include "llvm/Support/CommandLine.h" #include "llvm/TableGen/Error.h" @@ -31,13 +27,6 @@ // Utility Functions //===----------------------------------------------------------------------===// -std::string mlir::tblgen::getParameterAccessorName(StringRef name) { - assert(!name.empty() && "parameter has empty name"); - auto ret = "get" + name.str(); - ret[3] = llvm::toUpper(ret[3]); // uppercase first letter of the name - return ret; -} - /// Find all the AttrOrTypeDef for the specified dialect. If no dialect /// specified and can only find one dialect's defs, use that. static void collectAllDefs(StringRef selectedDialect, @@ -287,18 +276,20 @@ void DefGen::emitAccessors() { for (auto ¶m : params) { - Method *m = defCls.addMethod( - param.getCppAccessorType(), getParameterAccessorName(param.getName()), - def.genStorageClass() ? Method::Const : Method::ConstDeclaration); - // Generate accessor definitions only if we also generate the storage - // class. Otherwise, let the user define the exact accessor definition. - if (!def.genStorageClass()) - continue; - auto scope = m->body().indent().scope("return getImpl()->", ";"); - if (isa(param)) - m->body() << formatv("getType().cast<{0}>()", param.getCppType()); - else - m->body() << param.getName(); + for (StringRef accessorName : param.getAccessorNames()) { + Method *m = defCls.addMethod( + param.getCppAccessorType(), accessorName, + def.genStorageClass() ? Method::Const : Method::ConstDeclaration); + // Generate accessor definitions only if we also generate the storage + // class. Otherwise, let the user define the exact accessor definition. + if (!def.genStorageClass()) + continue; + auto scope = m->body().indent().scope("return getImpl()->", ";"); + if (isa(param)) + m->body() << formatv("getType().cast<{0}>()", param.getCppType()); + else + m->body() << param.getName(); + } } } diff --git a/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.h b/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.h --- a/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.h +++ b/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.h @@ -20,11 +20,6 @@ void generateAttrOrTypeFormat(const AttrOrTypeDef &def, MethodBody &parser, MethodBody &printer); -/// From the parameter name, get the name of the accessor function in camelcase. -/// The first letter of the parameter is upper-cased and prefixed with "get". -/// E.g. 'value' -> 'getValue'. -std::string getParameterAccessorName(llvm::StringRef name); - } // namespace tblgen } // namespace mlir diff --git a/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp b/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp --- a/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp +++ b/mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp @@ -56,9 +56,14 @@ /// Returns the name of the parameter. StringRef getName() const { return param.getName(); } + /// Get the accessor method name. + std::string getAccessorName() const { + return param.getAccessorNames().front(); + } + /// Generate the code to check whether the parameter should be printed. MethodBody &genPrintGuard(FmtContext &ctx, MethodBody &os) const { - std::string self = getParameterAccessorName(getName()) + "()"; + std::string self = getAccessorName() + "()"; ctx.withSelf(self); os << tgfmt("($_self", &ctx); if (llvm::Optional defaultValue = getParam().getDefaultValue()) { @@ -718,7 +723,7 @@ void DefFormat::genVariablePrinter(ParameterElement *el, FmtContext &ctx, MethodBody &os, bool skipGuard) { const AttrOrTypeParameter ¶m = el->getParam(); - ctx.withSelf(getParameterAccessorName(param.getName()) + "()"); + ctx.withSelf(param.getAccessorNames().front() + "()"); // Guard the printer on the presence of optional parameters and that they // aren't equal to their default values (if they have one). @@ -811,9 +816,7 @@ FormatElement *param = arg; if (auto *ref = dyn_cast(arg)) param = ref->getArg(); - os << ",\n" - << getParameterAccessorName(cast(param)->getName()) - << "()"; + os << ",\n" << cast(param)->getAccessorName() << "()"; } os.unindent() << ");\n"; }