Fix typos. NFC. (#20503)

Author: kuhar

compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.cpp

@@ -569,7 +569,7 @@ static Value foldTransferGatherFromStep(TransferGatherOp gatherOp) {
         int64_t resultDim = cast<AffineDimExpr>(map.getResult(0)).getPosition();
 
         // If the map is indexing along the memory dimension, and the vector is
-        // contigious, this is a contigious load on this dimension.
+        // contiguous, this is a contiguous load on this dimension.
         if (resultDim == index) {
           return {Value(), AffineMap(), true};
         }

compiler/src/iree/compiler/Codegen/Dialect/VectorExt/IR/VectorExtOps.td

@@ -139,22 +139,22 @@ def IREEVectorExt_TransferGatherOp : IREEVectorExt_PureOp<"transfer_gather", [
     the result vector and the memory dimensions being indexed.
 
     The operation is a generalization of `vector.transfer_read` op, where the
-    slice from which the read is performed is not guranteed to be contigious,
+    slice from which the read is performed is not guranteed to be contiguous,
     and instead how the slice is gathered is defined explicitly in the
     operation.
 
     The operation can be thought of as:
-      1. A contigious slice gathered from the source as described by the operation
-      2. A `vector.transfer_read` on the contigious slice
+      1. A contiguous slice gathered from the source as described by the operation
+      2. A `vector.transfer_read` on the contiguous slice
 
     The operation defines `permutation_map`, `padding`, `mask`, `in_bounds` in
     the same way as `vector.transfer_read` defines, but on the inferred
-    contigious slice.
+    contiguous slice.
 
-    The other parameters of the operation define how the contigious slice is
+    The other parameters of the operation define how the contiguous slice is
     gathered from the source. `indices` define a base to offset the source by.
     `indexed` defines for each dimension if the dimension is gathered or
-    contigious.
+    contiguous.
 
     The `indices` contains a base to offset the source by. The `indexed` array
     defines if a dimension is gathered or not. For example, for the following
@@ -194,7 +194,7 @@ def IREEVectorExt_TransferGatherOp : IREEVectorExt_PureOp<"transfer_gather", [
     ```
 
     With these additional parameters, the operation can define a supervector
-    read from a non-contigious slice. For example:
+    read from a non-contiguous slice. For example:
 
     ```
     source: memref<8192x8x16xf32>
@@ -222,7 +222,7 @@ def IREEVectorExt_TransferGatherOp : IREEVectorExt_PureOp<"transfer_gather", [
     The crucial structure of the operation relies on the index_vec and
     the result vector's indexing being defined based on the dimensions of the
     memory. This mapping can be exploited to simplify gathered dimensions
-    to contigious dimensions.
+    to contiguous dimensions.
   }];
 
   let extraClassDeclaration = [{

compiler/src/iree/compiler/Codegen/LLVMGPU/KernelConfig.cpp

@@ -96,7 +96,7 @@ llvm::cl::opt<bool>
                  llvm::cl::desc("force use of wmma operations for tensorcore"),
                  llvm::cl::init(false));
 
-/// Flag used to toggle using mma.sync vs wmma when targetting tensorcore.
+/// Flag used to toggle using mma.sync vs wmma when targeting tensorcore.
 llvm::cl::opt<bool>
     clGPUUseMMASync("iree-codegen-llvmgpu-use-mma-sync",
                     llvm::cl::desc("force use mma sync instead of wmma ops"),
@@ -160,7 +160,7 @@ static bool needsLoweringConfigPropagation(
 static SmallVector<TileWorkgroupSizePair>
 getMatmulConfig(IREE::GPU::TargetAttr target) {
   SmallVector<TileWorkgroupSizePair> tileSizes;
-  // Pick tile size so that M*K and K*N dividible by wgSize * \*vecSize=*\4.
+  // Pick tile size so that M*K and K*N divisible by wgSize * \*vecSize=*\4.
   // This way workgroup memory copy don't need to be masked. Once we support
   // masked load we can get performance out of more configuration.
 
@@ -189,7 +189,7 @@ getTensorCoreConfig(SmallVectorImpl<TileWorkgroupSizePair> &tileSizes,
                     Type elementType, int64_t M, int64_t N, int64_t K) {
   // Based on early analysis we found that 128x256x32_3 gives acceptable
   // performance across many of the large matrix sizes for f16 and fp32. This
-  // needs to be refined into a better strategy based on empircal data but this
+  // needs to be refined into a better strategy based on empirical data but this
   // gives us a quick solution to achieve performance in the right order of
   // magnitude for large square like cases.
   int64_t parallelDim = M * N;
@@ -1164,7 +1164,7 @@ static LogicalResult setAttentionIntrinsicBasedVectorDistributionConfig(
   // The subgroup distribution in attention is controlled by the second matmul
   // (Parallel dimension distribution is usually (almost always) controlled by
   // the last reduction operation in a dispatch). Since VectorDistribution
-  // doesn't have logic to set subgroup and thread layouts seperately, we
+  // doesn't have logic to set subgroup and thread layouts separately, we
   // explicitly set the subgroup count for the first matmul as well,
   // corresponding to what the second matmul dictates.
 
@@ -1624,7 +1624,7 @@ static LogicalResult setContractConfig(IREE::GPU::TargetAttr target,
 
   // Send very skinny, {2-4}xNxK and Mx{2-4}xK, matmuls to the vector reduction
   // pipeline, similar to matvec. Note: Because of reassociation in the vector
-  // reduction pipeline, this may lead to precission loss. If this ever becomes
+  // reduction pipeline, this may lead to precision loss. If this ever becomes
   // an issue, we can hide this behind a flag.
   if (llvm::all_equal({contractionDims->m.size(), contractionDims->n.size(),
                        contractionDims->k.size(), size_t{1}}) &&
@@ -2391,7 +2391,7 @@ static LogicalResult setTransposeConfig(mlir::FunctionOpInterface entryPoint,
 
   // Workgroup size contains 8 warps. Configured with 8 threads on fastest
   // moving dimension so each thread can execute a vectorized copy of 4
-  // contigious elements at a time from the 32 block.
+  // contiguous elements at a time from the 32 block.
   std::array<int64_t, 3> workgroupSize = {8, 32, 1};
 
   return setOpConfigAndEntryPointFnTranslation(
@@ -2470,7 +2470,7 @@ static LogicalResult setArgmaxUkernelConfig(
 }
 
 /// Decides the tiling and distribution parameters for one convolution
-/// dimension. Returns true if we can succesfully deduce.
+/// dimension. Returns true if we can successfully deduce.
 ///
 /// - `inputDim` is the size of the dimension to be distributed.
 /// - `residualThreads` is the remaining threads we can distribute.
@@ -2512,7 +2512,7 @@ static bool distributeToOneDim(const int64_t inputDim,
 
 /// Decides the tiling and distribution parameters for two convolution window
 /// dimensions to two workgroup dimensions as a square. Returns true if we can
-/// succesfully deduce.
+/// successfully deduce.
 static bool distributeToSquare(const int64_t oh, const int64_t ow,
                                int64_t &residualThreads,
                                int64_t &residualTilingFactor,

compiler/src/iree/compiler/Codegen/LLVMGPU/TransformExtensions/LLVMGPUExtensionsOps.td

@@ -341,7 +341,7 @@ def VectorToMMAConversionOp : Op<Transform_Dialect, "iree.vector.vector_to_mma_c
      ReportTrackingListenerFailuresOpTrait]> {
   let description = [{
     This converts slices of operations containing vector.contract op into
-    mma operations, targetting warp level tensorcore operations. If the vector
+    mma operations, targeting warp level tensorcore operations. If the vector
     operations are bigger than the native mma size it will first split up those
     vector operations.
 

compiler/src/iree/compiler/Codegen/SPIRV/KernelConfig.cpp

@@ -86,7 +86,7 @@ static bool fusedOpMayUseExtraSharedMemory(linalg::LinalgOp matmul) {
 //===----------------------------------------------------------------------===//
 
 /// Decides the tiling and distribution parameters for one convolution
-/// dimension. Returns true if we can succesfully deduce.
+/// dimension. Returns true if we can successfully deduce.
 ///
 /// - `inputDim` is the size of the dimension to be distributed.
 /// - `residualThreads` is the remaining threads we can distribute.
@@ -126,7 +126,7 @@ static bool tileConvOneDim(const int64_t inputDim, const bool isInnerMostDim,
 
 /// Decides the tiling and distribution parameters for two convolution window
 /// dimensions to two workgroup dimensions as a square. Returns true if we can
-/// succesfully deduce.
+/// successfully deduce.
 static bool tileConvSquare(const int64_t oh, const int64_t ow,
                            int64_t &residualThreads,
                            int64_t &residualTilingFactor,

compiler/src/iree/compiler/Codegen/SPIRV/SPIRVVectorizeLoadStore.cpp

@@ -767,7 +767,7 @@ static Value predicateMaybeMaskedScalarTransfer(
 }
 
 /// Scalarizes remaining vector transfer that couldn't be converted to
-/// vevtor load operations.
+/// vector load operations.
 
 /// This is very specific to SPIR-V as pointer cannot be casted to vector type
 /// if any of the memory access is not vector.