[mlir] Fixed a number of typos

[GabrielRavier]

Summary

I went over the output of the following mess of a command:

(ulimit -m 2000000; ulimit -v 2000000; git ls-files -z | parallel --xargs -0 cat | aspell list --mode=none --ignore-case | grep -E '^[A-Za-z][a-z]*$' | sort | uniq -c | sort -n | grep -vE '.{25}' | aspell pipe -W3 | grep : | cut -d' ' -f2 | less)

and proceeded to spend a few days looking at it to find probable typos
and fixed a few hundred of them in all of the llvm project (note, the
ones I found are not anywhere near all of them, but it seems like a
good start).

[dcaballe]

Wow, thank you so much for doing this! LGTM.

[aartbik]

Thanks for the PyTACO fix

mlir/docs/Bindings/Python.md

 
 -   a generic `Value` object; or
 -   a concrete `BlockArgument` or `OpResult` object.
 
 The former provides all the generic functionality such as comparison, type
 access and printing. The latter provide access to the defining block or
 operation and the position of the value within it. By default, the generic
 `Value` objects are returned from IR traversals. Downcasting is implemented
-through concrete subclass constructors, similarly to attribtues and types:
+through concrete subclass constructors, similarly to attributes and types:
 
 ```python
 from mlir.ir import BlockArgument, OpResult, Value
 
 value = ...
 
 # Set `concrete` to the specific value subclass.
 try:

mlir/docs/TargetLLVMIR.md

 }
 ```
 
 #### Bare Pointer Calling Convention For Unranked MemRef
 
 The "bare pointer" calling convention does not support unranked memrefs as their
 shape cannot be known at compile time.
 
-### Generic alloction and deallocation functions
+### Generic allocation and deallocation functions
 
 When converting the Memref dialect, allocations and deallocations are converted
 into calls to `malloc` (`aligned_alloc` if aligned allocations are requested)
 and `free`. However, it is possible to convert them to more generic functions
 which can be implemented by a runtime library, thus allowing custom allocation
 strategies or runtime profiling. When the conversion pass is  instructed to
 perform such operation, the names of the calles are
 `_mlir_memref_to_llvm_alloc`, `_mlir_memref_to_llvm_aligned_alloc` and

mlir/include/mlir/Analysis/DataFlowFramework.h

public:
   /// calling `addDependency` between the input state and a program point,
   /// indicating that, if the state is updated, the solver should invoke `solve`
   /// on the program point. The dependent point does not have to be the same as
   /// the provided point. An update to a state is propagated by calling
   /// `propagateIfChange` on the state. If the state has changed, then all its
   /// dependents are placed on the worklist.
   ///
   /// The dependency graph does not need to be static. Each invocation of
-  /// `visit` can add new dependencies, but these dependecies will not be
+  /// `visit` can add new dependencies, but these dependencies will not be
   /// dynamically added to the worklist because the solver doesn't know what
   /// will provide a value for then.
   virtual LogicalResult visit(ProgramPoint point) = 0;
 
 protected:
   /// Create a dependency between the given analysis state and program point
   /// on this analysis.
   void addDependency(AnalysisState *state, ProgramPoint point);

mlir/include/mlir/Analysis/Presburger/PWMAFunction.h

public:
   /// corresponding to the added variables are initialized to zero. Return the
   /// absolute column position (i.e., not relative to the kind of variable)
   /// of the first added variable.
   unsigned insertVar(VarKind kind, unsigned pos, unsigned num = 1);
 
   /// Remove the specified range of vars.
   void removeVarRange(VarKind kind, unsigned varStart, unsigned varLimit);
 
-  /// Given a MAF `other`, merges local variables such that both funcitons
+  /// Given a MAF `other`, merges local variables such that both functions
   /// have union of local vars, without changing the set of points in domain or
   /// the output.
   void mergeLocalVars(MultiAffineFunction &other);
 
   /// Return whether the outputs of `this` and `other` agree wherever both
   /// functions are defined, i.e., the outputs should be equal for all points in
   /// the intersection of the domains.
   bool isEqualWhereDomainsOverlap(MultiAffineFunction other) const;

mlir/include/mlir/Analysis/Presburger/Simplex.h

 /// LexSimplex::findIntegerLexMin simultaneously for every value of the symbols
 /// in the domain.
 ///
 /// Often, the pivot to be performed is the same for all values of the symbols,
 /// in which case we just do it. For example, if the symbolic sample of a row is
 /// negative for all values in the symbol domain, the row needs to be pivoted
 /// irrespective of the precise value of the symbols. To answer queries like
 /// "Is this symbolic sample always negative in the symbol domain?", we maintain
-/// a `LexSimplex domainSimplex` correponding to the symbol domain.
+/// a `LexSimplex domainSimplex` corresponding to the symbol domain.
 ///
 /// In other cases, it may be that the symbolic sample is violated at some
 /// values in the symbol domain and not violated at others. In this case,
 /// the pivot to be performed does depend on the value of the symbols. We
 /// handle this by splitting the symbol domain. We run the algorithm for the
 /// case where the row isn't violated, and then come back and run the case
 /// where it is.
 class SymbolicLexSimplex : public LexSimplexBase {

mlir/include/mlir/Bindings/Python/Attributes.td

 
 // A mapping between the attribute storage type and the corresponding Python
 // type. There is not necessarily a 1-1 match for non-builtin attributes.
 class PythonAttr<string c, string p> {
   string cppStorageType = c;
   string pythonType = p;
 }
 
-// Mappings between supported builtin attribtues and Python types.
+// Mappings between supported builtin attributes and Python types.
 def : PythonAttr<"::mlir::Attribute", "_ods_ir.Attribute">;
 def : PythonAttr<"::mlir::BoolAttr", "_ods_ir.BoolAttr">;
 def : PythonAttr<"::mlir::IntegerAttr", "_ods_ir.IntegerAttr">;
 def : PythonAttr<"::mlir::FloatAttr", "_ods_ir.FloatAttr">;
 def : PythonAttr<"::mlir::StringAttr", "_ods_ir.StringAttr">;
 def : PythonAttr<"::mlir::DenseElementsAttr", "_ods_ir.DenseElementsAttr">;
 def : PythonAttr<"::mlir::DenseIntElementsAttr", "_ods_ir.DenseIntElementsAttr">;
 def : PythonAttr<"::mlir::DenseFPElementsAttr", "_ods_ir.DenseFPElementsAttr">;
 
 #endif

mlir/include/mlir/Dialect/AMDGPU/AMDGPU.td

Results<(outs AnyTypeOf<[BF16, F16, F32, I32, I8,
                               VectorOfLengthAndType<[2, 4, 8], [F16, BF16]>,
                               VectorOfLengthAndType<[2, 4, 8, 16], [I8]>]>:$value)> {
 
   let summary = "Raw Buffer load, exposing GCN features";
   let description = [{
     The `amdgpu.raw_buffer_load` op is a wrapper around the buffer load intrinsics
     available on AMD GPUs, including extensions in newer GPUs.
 
-    The index into the buffer is computed as for `memref.load` with the additon
+    The index into the buffer is computed as for `memref.load` with the addition
     of `indexOffset` and `sgprOffset` (which **may or may not** be considered
     in bounds checks and includes any offset present on the memref type if it's
     non-zero).
 
     All indices and offsets are in units of the memref's data type and are
     converted to bytes during lowering.
 
     When a load is out of bounds, the instruction returns zero.

mlir/include/mlir/Dialect/GPU/IR/GPUBase.td

 
 def GPU_AsyncOpInterface : OpInterface<"AsyncOpInterface"> {
   let description = [{
     Interface for GPU operations that execute asynchronously on the device.
 
     GPU operations implementing this interface take a list of dependencies
     as `gpu.async.token` arguments and optionally return a `gpu.async.token`.
 
-    The op doesn't start executing until all depent ops producing the async
+    The op doesn't start executing until all dependent ops producing the async
     dependency tokens have finished executing.
 
     If the op returns a token, the op merely schedules the execution on the
     device and returns immediately, without waiting for the execution to
     complete. On the hand, if the op does not return a token, the op will wait
     for the execution to complete.
   }];
   let cppNamespace = "::mlir::gpu";

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

def NVVM_MmaOp : NVVM_Op<"mma.sync", [AttrSizedOperandSegments]> {
 
   let description = [{
     The `nvvm.mma.sync` operation collectively performs the operation
     `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
+    instruction shapes given as "mMnNkK". The below table describes the possibilities
     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 
-    `multiplicandDataType`. It specifies how the multiply-and-acumulate is
+    `multiplicandDataType`. It specifies how the multiply-and-accumulate is
     performed and is either `xor_popc` or `and_poc`. The default is `xor_popc`.
 
     `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 
     which case overflow wraps from one end of the range to the other.

mlir/include/mlir/Dialect/NVGPU/IR/NVGPU.td

 //===-- NVGPU.td - NVGPU dialect operation definitions *- tablegen -*------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the basic operations for the NVGPU dialect.
 //
 // This NVGPU provides a bridge between the target agnostic GPU and Vector
 // dialects and lower level NVVM dialect. This allow representing PTX specific
 // operations while using MLIR high level concepts like memref and 2-D vector.
 //
-// Ops semantic are going to be based on vendor specific PTX defintion:
+// Ops semantic are going to be based on vendor specific PTX definition:
 // https://docs.nvidia.com/cuda/parallel-thread-execution/index.html
 //
 //===----------------------------------------------------------------------===//
 
 #ifndef NVGPU
 #define NVGPU
 
 include "mlir/Interfaces/SideEffectInterfaces.td"

mlir/include/mlir/Dialect/Vector/Transforms/VectorDistribution.h

 class RewritePatternSet;
 namespace vector {
 
 struct WarpExecuteOnLane0LoweringOptions {
   /// Lamdba function to let users allocate memory needed for the lowering of
   /// WarpExecuteOnLane0Op.
   /// The function needs to return an allocation that the lowering can use as
   /// temporary memory. The allocation needs to match the shape of the type (the
-  /// type may be VectorType or a scalar) and be availble for the current warp.
+  /// type may be VectorType or a scalar) and be available for the current warp.
   /// If there are several warps running in parallel the allocation needs to be
   /// split so that each warp has its own allocation.
   using WarpAllocationFn =
       std::function<Value(Location, OpBuilder &, WarpExecuteOnLane0Op, Type)>;
   WarpAllocationFn warpAllocationFn = nullptr;
 
   /// Lamdba function to let user emit operation to syncronize all the thread
   /// within a warp. After this operation all the threads can see any memory

mlir/include/mlir/IR/BuiltinAttributes.td

def Builtin_SymbolRefAttr : Builtin_Attr<"SymbolRef", [
   let extraClassDeclaration = [{
     static SymbolRefAttr get(MLIRContext *ctx, StringRef value,
                              ArrayRef<FlatSymbolRefAttr> nestedRefs);
     /// Convenience getters for building a SymbolRefAttr with no path, which is
     /// known to produce a FlatSymbolRefAttr.
     static FlatSymbolRefAttr get(StringAttr value);
     static FlatSymbolRefAttr get(MLIRContext *ctx, StringRef value);
 
-    /// Convenience getter for buliding a SymbolRefAttr based on an operation
+    /// Convenience getter for building a SymbolRefAttr based on an operation
     /// that implements the SymbolTrait.
     static FlatSymbolRefAttr get(Operation *symbol);
 
     /// Returns the name of the fully resolved symbol, i.e. the leaf of the
     /// reference path.
     StringAttr getLeafReference() const;
   }];
   let skipDefaultBuilders = 1;

mlir/include/mlir/Transforms/Passes.td

let description = [{
     This pass implements control-flow sink on operations that implement
     `RegionBranchOpInterface` by moving dominating operations whose only uses
     are in a conditionally-executed regions into those regions so that
     executions paths where their results are not needed do not perform
     unnecessary computations.
 
     This is similar (but opposite) to loop-invariant code motion, which hoists
     operations out of regions executed more than once. The implementation of
-    control-flow sink uses a simple and conversative cost model: operations are
+    control-flow sink uses a simple and conservative cost model: operations are
     never duplicated and are only moved into singly-executed regions.
 
     It is recommended to run canonicalization first to remove unreachable
     blocks: ops in unreachable blocks may prevent other operations from being
     sunk as they may contain uses of their results
   }];
   let constructor = "::mlir::createControlFlowSinkPass()";
   let statistics = [

mlir/lib/Conversion/GPUCommon/GPUOpsLowering.cpp

auto globalType =
       LLVM::LLVMArrayType::get(llvmI8, formatString.size_in_bytes());
   LLVM::GlobalOp global;
   {
     ConversionPatternRewriter::InsertionGuard guard(rewriter);
     rewriter.setInsertionPointToStart(moduleOp.getBody());
     global = rewriter.create<LLVM::GlobalOp>(
         loc, globalType,
         /*isConstant=*/true, LLVM::Linkage::Internal, stringConstName,
-        rewriter.getStringAttr(formatString), /*allignment=*/0, addressSpace);
+        rewriter.getStringAttr(formatString), /*alignment=*/0, addressSpace);
   }
 
   // Get a pointer to the format string's first element
   Value globalPtr = rewriter.create<LLVM::AddressOfOp>(loc, global);
   Value stringStart = rewriter.create<LLVM::GEPOp>(
       loc, i8Ptr, globalPtr, ArrayRef<LLVM::GEPArg>{0, 0});
 
   // Construct arguments and function call

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

bool shouldBeDetensored(Operation *op, TypeConverter typeConverter) {
   GenericOp genericOp = dyn_cast_or_null<GenericOp>(op);
   return genericOp &&
          llvm::all_of(
              genericOp.getInputAndOutputOperands(), [&](OpOperand *opOperand) {
                return !typeConverter.isLegal(opOperand->get().getType());
              });
 }
 
-/// A conversion patttern for detensoring `linalg.generic` ops.
+/// A conversion pattern for detensoring `linalg.generic` ops.
 class DetensorizeGenericOp : public OpConversionPattern<GenericOp> {
 public:
   using OpConversionPattern::OpConversionPattern;
   LogicalResult
   matchAndRewrite(GenericOp op, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override {
     Block *originalBlock = op->getBlock();
 
void compute(FunctionOpInterface func,
       // ops should not be detensored as well.
 
       DenseSet<BlockArgument> blockArgsToRemove;
 
       for (auto &blockArg : blockArgsToDetensor) {
         Block *block = blockArg.getParentBlock();
 
         // For the potentially detensorable block argument, find the
-        // correpsonding operands in predecessor blocks.
+        // corresponding operands in predecessor blocks.
         for (PredecessorIterator pred = block->pred_begin();
              pred != block->pred_end(); ++pred) {
           BranchOpInterface terminator =
               dyn_cast<BranchOpInterface>((*pred)->getTerminator());
           auto blockOperands =
               terminator.getSuccessorOperands(pred.getSuccessorIndex());
 
           if (blockOperands.empty() ||

mlir/lib/Dialect/Math/Transforms/PolynomialApproximation.cpp

Value zeroMask = builder.create<arith::CmpFOp>(arith::CmpFPredicate::OEQ,
                                                  op.getOperand(), cstZero);
   Value posInfMask = builder.create<arith::CmpFOp>(arith::CmpFPredicate::OEQ,
                                                    op.getOperand(), cstPosInf);
 
   // Filter out invalid values:
   //  • x == 0     -> -INF
   //  • x < 0      ->  NAN
   //  • x == +INF  -> +INF
-  Value aproximation = builder.create<arith::SelectOp>(
+  Value approximation = builder.create<arith::SelectOp>(
       zeroMask, cstMinusInf,
       builder.create<arith::SelectOp>(
           invalidMask, cstNan,
           builder.create<arith::SelectOp>(posInfMask, cstPosInf, x)));
 
-  rewriter.replaceOp(op, aproximation);
+  rewriter.replaceOp(op, approximation);
 
   return success();
 }
 
 namespace {
 struct LogApproximation : public LogApproximationBase<math::LogOp> {
   using LogApproximationBase::LogApproximationBase;
 

mlir/lib/Dialect/Tosa/Transforms/TosaInferShapes.cpp

 public:
   void runOnOperation() override {
     func::FuncOp func = getOperation();
 
     IRRewriter rewriter(func.getContext());
 
     propagateShapesInRegion(func.getBody());
 
-    // Insert UnrealizedConversionCasts to guarantee ReturnOp agress with
+    // Insert UnrealizedConversionCasts to guarantee ReturnOp agrees with
     // the FuncOp type.
     func.walk([&](func::ReturnOp op) {
       func::FuncOp parent = dyn_cast<func::FuncOp>(op->getParentOp());
       if (!parent)
         return;
 
       rewriter.setInsertionPoint(op);
       FunctionType funcTy = func.getFunctionType();

mlir/lib/ExecutionEngine/AsyncRuntime.cpp

extern "C" int64_t mlirAsyncRuntimeAddTokenToGroup(AsyncToken *token,
 
   if (State(token->state).isAvailableOrError()) {
     // Update group pending tokens immediately and maybe run awaiters.
     onTokenReady();
 
   } else {
     // Update group pending tokens when token will become ready. Because this
     // will happen asynchronously we must ensure that `group` is alive until
-    // then, and re-ackquire the lock.
+    // then, and re-acquire the lock.
     group->addRef();
 
     token->awaiters.emplace_back([group, onTokenReady]() {
       // Make sure that `dropRef` does not destroy the mutex owned by the lock.
       {
         std::unique_lock<std::mutex> lockGroup(group->mu);
         onTokenReady();
       }

mlir/lib/Target/Cpp/TranslateToCpp.cpp

struct CppEmitter {
   /// - emits nothing if no value produced by op;
   /// Emits final '=' operator where a type is produced. Returns failure if
   /// any result type could not be converted.
   LogicalResult emitAssignPrefix(Operation &op);
 
   /// Emits a label for the block.
   LogicalResult emitLabel(Block &block);
 
-  /// Emits the operands and atttributes of the operation. All operands are
+  /// Emits the operands and attributes of the operation. All operands are
   /// emitted first and then all attributes in alphabetical order.
   LogicalResult emitOperandsAndAttributes(Operation &op,
                                           ArrayRef<StringRef> exclude = {});
 
   /// Emits the operands of the operation. All operands are emitted in order.
   LogicalResult emitOperands(Operation &op);
 
   /// Return the existing or a new name for a Value.

mlir/lib/Target/LLVMIR/Dialect/OpenACC/OpenACCToLLVMIRTranslation.cpp

if (failed(processOperands(builder, moduleTranslation, op, op.copyOperands(),
                              names, index, mapperAllocas)))
     return failure();
 
   if (failed(processOperands(
           builder, moduleTranslation, op, op.copyinOperands(), totalNbOperand,
           kDeviceCopyinFlag | kHoldFlag, flags, names, index, mapperAllocas)))
     return failure();
 
-  // TODO copyin readonly currenlty handled as copyin. Update when extension
+  // TODO copyin readonly currently handled as copyin. Update when extension
   // available.
   if (failed(processOperands(builder, moduleTranslation, op,
                              op.copyinReadonlyOperands(), totalNbOperand,
                              kDeviceCopyinFlag | kHoldFlag, flags, names, index,
                              mapperAllocas)))
     return failure();
 
   if (failed(processOperands(
           builder, moduleTranslation, op, op.copyoutOperands(), totalNbOperand,
           kHostCopyoutFlag | kHoldFlag, flags, names, index, mapperAllocas)))
     return failure();
 
-  // TODO copyout zero currenlty handled as copyout. Update when extension
+  // TODO copyout zero currently handled as copyout. Update when extension
   // available.
   if (failed(processOperands(builder, moduleTranslation, op,
                              op.copyoutZeroOperands(), totalNbOperand,
                              kHostCopyoutFlag | kHoldFlag, flags, names, index,
                              mapperAllocas)))
     return failure();
 
   if (failed(processOperands(
           builder, moduleTranslation, op, op.createOperands(), totalNbOperand,
           kCreateFlag | kHoldFlag, flags, names, index, mapperAllocas)))
     return failure();
 
-  // TODO create zero currenlty handled as create. Update when extension
+  // TODO create zero currently handled as create. Update when extension
   // available.
   if (failed(processOperands(builder, moduleTranslation, op,
                              op.createZeroOperands(), totalNbOperand,
                              kCreateFlag | kHoldFlag, flags, names, index,
                              mapperAllocas)))
     return failure();
 
   if (failed(processOperands(

mlir/lib/Transforms/Utils/RegionUtils.cpp

for (int operand : llvm::seq<int>(0, lhsIt->getNumOperands())) {
       // Otherwise, these operands must have the same logical order within the
       // parent block.
       if (leaderData.getOrderOf(lhsOperand) != blockData.getOrderOf(rhsOperand))
         return failure();
     }
 
     // If the lhs or rhs has external uses, the blocks cannot be merged as the
     // merged version of this operation will not be either the lhs or rhs
-    // alone (thus semantically incorrect), but some mix dependending on which
+    // alone (thus semantically incorrect), but some mix depending on which
     // block preceeded this.
     // TODO allow merging of operations when one block does not dominate the
     // other
     if (rhsIt->isUsedOutsideOfBlock(mergeBlock) ||
         lhsIt->isUsedOutsideOfBlock(leaderBlock)) {
       return failure();
     }
   }

mlir/python/mlir/dialects/linalg/opdsl/lang/affine.py

class AffineBuildState:
 
   def get_dim(self, dimname: str) -> int:
     """Gets the dim position given a name."""
     pos = self.all_dims.get(dimname)
     if pos is None:
       if not self.allow_new_dims:
         raise ValueError(
             f"New dimensions not allowed in the current affine expression: "
-            f"Requested '{dimname}', Availble: {self.all_dims}")
+            f"Requested '{dimname}', Available: {self.all_dims}")
       pos = len(self.all_dims)
       self.all_dims[dimname] = pos
     self.local_dims[dimname] = pos
     return pos
 
   def get_symbol(self, symname: str) -> int:
     """Geta a symbol position given a name."""
     pos = self.all_symbols.get(symname)
     if pos is None:
       if not self.allow_new_symbols:
         raise ValueError(
             f"New symbols not allowed in the current affine expression: "
-            f"Requested '{symname}', Availble: {self.all_symbols}")
+            f"Requested '{symname}', Available: {self.all_symbols}")
       pos = len(self.all_symbols)
       self.all_symbols[symname] = pos
     self.local_symbols[symname] = pos
     return pos
 
   @property
   def local_dim_count(self) -> int:
     return len(self.local_dims)

mlir/test/Conversion/NVGPUToNVVM/nvgpu-to-nvvm.mlir

func.func @m16n8k16_fp16(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf16>) -> vector<2x2xf16> {
   // CHECK: llvm.mlir.undef : !llvm.array<2 x vector<2xf16>>
   // CHECK-DAG: llvm.insertvalue {{%.+}}, {{%.+}}[0] : !llvm.array<2 x vector<2xf16>>
   // CHECK-DAG: llvm.insertvalue {{%.+}}, {{%.+}}[1] : !llvm.array<2 x vector<2xf16>>      
   return %d : vector<2x2xf16>
 }
 
 // -----
 
-// Same as above but with fp32 acumulation type.
+// Same as above but with fp32 accumulation type.
 
 // CHECK-LABEL: @m16n8k16_fp16_fp32
 func.func @m16n8k16_fp16_fp32(%arg0: vector<4x2xf16>, %arg1: vector<2x2xf16>, %arg2: vector<2x2xf32>) -> vector<2x2xf32> {
   // We just need to check the mma instruction and the manipulatin of the result.
   // CHECK: [[d:%.+]] = nvvm.mma.sync
   // CHECK-SAME: shape = #nvvm.shape<m = 16, n = 8, k = 16>
   // CHECK-SAME: (vector<2xf16>, vector<2xf16>, f32) -> !llvm.struct<(f32, f32, f32, f32)>
   %d = nvgpu.mma.sync (%arg0, %arg1, %arg2) {mmaShape = [16, 8, 16]} : (vector<4x2xf16>, vector<2x2xf16>, vector<2x2xf32>) -> vector<2x2xf32>    

mlir/test/Integration/Dialect/SparseTensor/taco/tools/mlir_pytaco.py

class _Stats:
 
   def get_formats(self, idx: int) -> Tuple[ModeFormat]:
     """Gets the ModeFormats for the ith temporary tensor."""
     return tuple(self._get_element(idx).dst_format.format_pack.formats)
 
 
 class _SparseValueInfo(enum.Enum):
   """Describes how a sparse tensor value is stored.
-  _UNPACKED: The sparse tensor value is stored as (coordnates, values) in
+  _UNPACKED: The sparse tensor value is stored as (coordinates, values) in
     Python.
   _PACKED: The sparse tensor value is stored as a C pointer to a packed MLIR
     sparse tensor.
   """
   _UNPACKED = 0
   _PACKED = 1
 
 
class Tensor:
 
   def is_dense(self) -> bool:
     """Returns true if the tensor doesn't have sparsity annotation."""
     return self.order == 0 or self._format is None
 
   def to_array(self) -> np.ndarray:
     """Returns the numpy array for the Tensor.
 
-    This is currenly only implemented for dense Tensor.
+    This is currently only implemented for dense Tensor.
     """
     if not self.is_dense():
       raise ValueError("Conversion from non-dense Tensor "
                        "to numpy array not supported yet.")
 
     self._sync_value()
 
     return self._dense_storage

mlir/tools/mlir-linalg-ods-gen/mlir-linalg-ods-yaml-gen.cpp

 
 // Template for Linalg named ops' ODS definitions. Parameters:
 // {0}: ODS/C++ op name
 // {1}: assembly op mnemonic
 // {2}: op interface list
 // {3}: documentation (summary + description)
 // {4}: op attribute list
 // {5}: builder methods taking standalone attribute parameters
-// {6}: additional method defintions
+// {6}: additional method definitions
 // {7}: additional methods for attributes used by indexing maps
 static const char structuredOpOdsHeaderFormat[] = R"FMT(
 //===----------------------------------------------------------------------===//
 // Op definition for {0}
 //===----------------------------------------------------------------------===//
 
 def {0} : LinalgStructuredBase_Op<"{1}", !listconcat([AttrSizedOperandSegments],
   /*extraInterfaces=*/[{2}])> {

mlir/tools/mlir-tblgen/AttrOrTypeFormatGen.cpp

if (hasOptional) {
     os << "do {\n";
     os.indent();
   }
 
   ArrayRef<ParameterElement *> params = el->getParams();
   using IteratorT = ParameterElement *const *;
   IteratorT it = params.begin();
 
-  // Find the last required parameter. Commas become optional aftewards.
+  // Find the last required parameter. Commas become optional afterwards.
   // Note: IteratorT's copy assignment is deleted.
   ParameterElement *lastReq = nullptr;
   for (ParameterElement *param : params)
     if (!param->isOptional())
       lastReq = param;
   IteratorT lastReqIt = lastReq ? llvm::find(params, lastReq) : params.begin();
 
   auto eachFn = [&](ParameterElement *el) { genVariableParser(el, ctx, os); };

mlir/tools/mlir-tblgen/OpFormatGen.cpp

private:
   /// The argument that is used to format the directive.
   FormatElement *arg;
 
   bool shouldBeQualifiedFlag = false;
 };
 
 /// This class represents a group of order-independent optional clauses. Each
 /// clause starts with a literal element and has a coressponding parsing
-/// element. A parsing element is a continous sequence of format elements.
+/// element. A parsing element is a continuous sequence of format elements.
 /// Each clause can appear 0 or 1 time.
 class OIListElement : public DirectiveElementBase<DirectiveElement::OIList> {
 public:
   OIListElement(std::vector<FormatElement *> &&literalElements,
                 std::vector<std::vector<FormatElement *>> &&parsingElements)
       : literalElements(std::move(literalElements)),
         parsingElements(std::move(parsingElements)) {}