[MLIR][arith] fix references to std.constant in comments

[Mogball]

Ensure verify is called in CAPI tests

[Mogball]

Fix ArithmeticPasses documentation

flang/include/flang/Optimizer/Dialect/FIROps.td

 
 def fir_HasValueOp : fir_Op<"has_value", [Terminator, HasParent<"GlobalOp">]> {
   let summary = "terminator for GlobalOp";
   let description = [{
     The terminator for a GlobalOp with a body.
 
     ```mlir
       global @variable : tuple<i32, f32> {
-        %0 = constant 45 : i32
-        %1 = constant 100.0 : f32
+        %0 = arith.constant 45 : i32
+        %1 = arith.constant 100.0 : f32
         %2 = fir.undefined tuple<i32, f32>
-        %3 = constant 0 : index
+        %3 = arith.constant 0 : index
         %4 = fir.insert_value %2, %0, %3 : (tuple<i32, f32>, i32, index) -> tuple<i32, f32>
-        %5 = constant 1 : index
+        %5 = arith.constant 1 : index
         %6 = fir.insert_value %4, %1, %5 : (tuple<i32, f32>, f32, index) -> tuple<i32, f32>
         fir.has_value %6 : tuple<i32, f32>
       }
     ```
   }];
 
   let arguments = (ins AnyType:$resval);
 
def fir_EmboxOp : fir_Op<"embox", [NoSideEffect, AttrSizedOperandSegments]> {
 
   let description = [{
     Create a boxed reference value. In Fortran, the implementation can require
     extra information about an entity, such as its type, rank, etc.  This
     auxilliary information is packaged and abstracted as a value with box type
     by the calling routine. (In Fortran, these are called descriptors.)
 
     ```mlir
-      %c1 = constant 1 : index
-      %c10 = constant 10 : index
+      %c1 = arith.constant 1 : index
+      %c10 = arith.constant 10 : index
       %5 = ... : !fir.ref<!fir.array<10 x i32>>
       %6 = fir.embox %5 : (!fir.ref<!fir.array<10 x i32>>) -> !fir.box<!fir.array<10 x i32>>
     ```
 
     The descriptor tuple may contain additional implementation-specific
     information through the use of additional attributes.
     Specifically,
         - shape: emboxing an array may require shape information (an array's
let description = [{
     a variable of type CHARACTER has both its data reference as well as a
     LEN type parameter.
 
     ```fortran
       CHARACTER(LEN=10) :: var
     ```
     ```mlir
       %4 = ...         : !fir.ref<!fir.array<10 x !fir.char<1>>>
-      %5 = constant 10 : i32
+      %5 = arith.constant 10 : i32
       %6 = fir.emboxchar %4, %5 : (!fir.ref<!fir.array<10 x !fir.char<1>>>, i32) -> !fir.boxchar<1>
     ```
 
     In the above `%4` is a memory reference to a buffer of 10 CHARACTER units.
     This buffer and its LEN value (10) are wrapped into a pair in `%6`.
   }];
 
   let arguments = (ins AnyReferenceLike:$memref, AnyIntegerLike:$len);
def fir_BoxDimsOp : fir_Op<"box_dims", [NoSideEffect]> {
 
   let description = [{
     Returns the triple of lower bound, extent, and stride for `dim` dimension
     of `val`, which must have a `box` type.  The dimensions are enumerated from
     left to right from 0 to rank-1. This operation has undefined behavior if
     `dim` is out of bounds.
 
     ```mlir
-      %c1   = constant 0 : i32
+      %c1   = arith.constant 0 : i32
       %52:3 = fir.box_dims %40, %c1 : (!fir.box<!fir.array<*:f64>>, i32) -> (index, index, index)
     ```
 
     The above is a request to return the left most row (at index 0) triple from
     the box. The triple will be the lower bound, extent, and byte-stride, which
     are the values encoded in a standard descriptor.
   }];
 
def fir_BoxIsArrayOp : fir_SimpleOp<"box_isarray", [NoSideEffect]> {
 
   let description = [{
     Determine if the boxed value has a positive (> 0) rank. This will return
     true if the originating box value was from a fir.embox with a memory
     reference value that had the type !fir.array<T> and/or a shape argument.
 
     ```mlir
       %r = ... : !fir.ref<i64>
-      %c_100 = constant 100 : index
+      %c_100 = arith.constant 100 : index
       %d = fir.shape %c_100 : (index) -> !fir.shape<1>
       %b = fir.embox %r(%d) : (!fir.ref<i64>, !fir.shape<1>) -> !fir.box<i64>
       %a = fir.box_isarray %b : (!fir.box<i64>) -> i1  // true
     ```
   }];
 
   let arguments = (ins fir_BoxType:$val);
 
let description = [{
 
     Note that the entity ssa-value must be of compile-time known size in order
     to use this operation.
 
     ```mlir
       %a = ... : !fir.array<10xtuple<i32, f32>>
       %f = ... : f32
       %o = ... : i32
-      %c = constant 1 : i32
+      %c = arith.constant 1 : i32
       %b = fir.insert_value %a, %f, %o, %c : (!fir.array<10x20xtuple<i32, f32>>, f32, i32, i32) -> !fir.array<10x20xtuple<i32, f32>>
     ```
   }];
 
   let arguments = (ins AnyCompositeLike:$adt, AnyType:$val, ArrayAttr:$coor);
   let results = (outs AnyCompositeLike);
 
   let assemblyFormat = [{
let description = [{
     Insert copies of a value into an entity with an array type.
     Returns a new ssa value with the same type as the original entity.
     The values are inserted at a contiguous range of indices in Fortran
     row-to-column element order as specified by lower and upper bound
     coordinates.
 
     ```mlir
       %a = fir.undefined !fir.array<10x10xf32>
-      %c = constant 3.0 : f32
+      %c = arith.constant 3.0 : f32
       %1 = fir.insert_on_range %a, %c, [0 : index, 7 : index, 0 : index, 2 : index] : (!fir.array<10x10xf32>, f32) -> !fir.array<10x10xf32>
     ```
 
     The first 28 elements of %1, with coordinates from (0,0) to (7,2), have
     the value 3.0.
   }];
 
   let arguments = (ins fir_SequenceType:$seq, AnyType:$val, ArrayAttr:$coor);

 def fir_DoLoopOp : region_Op<"do_loop",
     [DeclareOpInterfaceMethods<LoopLikeOpInterface>]> {
   let summary = "generalized loop operation";
   let description = [{
     Generalized high-level looping construct. This operation is similar to
     MLIR's `scf.for`.
 
     ```mlir
-      %l = constant 0 : index
-      %u = constant 9 : index
-      %s = constant 1 : index
+      %l = arith.constant 0 : index
+      %u = arith.constant 9 : index
+      %s = arith.constant 1 : index
       fir.do_loop %i = %l to %u step %s unordered {
         %x = fir.convert %i : (index) -> i32
         %v = fir.call @compute(%x) : (i32) -> f32
         %p = fir.coordinate_of %A, %i : (!fir.ref<!fir.array<?xf32>>, index) -> !fir.ref<f32>
         fir.store %v to %p : !fir.ref<f32>
       }
     ```
 
let description = [{
     A global variable or constant with initial values.
 
     The example creates a global variable (writable) named
     `@_QV_Mquark_Vvarble` with some initial values. The initializer should
     conform to the variable's type.
 
     ```mlir
       fir.global @_QV_Mquark_Vvarble : tuple<i32, f32> {
-        %1 = constant 1 : i32
-        %2 = constant 2.0 : f32
+        %1 = arith.constant 1 : i32
+        %2 = arith.constant 2.0 : f32
         %3 = fir.undefined tuple<i32, f32>
-        %z = constant 0 : index
-        %o = constant 1 : index
+        %z = arith.constant 0 : index
+        %o = arith.constant 1 : index
         %4 = fir.insert_value %3, %1, %z : (tuple<i32, f32>, i32, index) -> tuple<i32, f32>
         %5 = fir.insert_value %4, %2, %o : (tuple<i32, f32>, f32, index) -> tuple<i32, f32>
         fir.has_value %5 : tuple<i32, f32>
       }
     ```
   }];
 
   let arguments = (ins

mlir/docs/Bufferization.md

-   This is an example of a pass that is not split along dialect
         subdivisions.
 
 -   `tensor-constant-bufferize`
     ([code](https://github.com/llvm/llvm-project/blob/bc8acf2ce8ad6e8c9b1d97b2e02d3f4ad26e1d9d/mlir/lib/Dialect/StandardOps/Transforms/TensorConstantBufferize.cpp#L1),
     [test](https://github.com/llvm/llvm-project/blob/bc8acf2ce8ad6e8c9b1d97b2e02d3f4ad26e1d9d/mlir/test/Dialect/Standard/tensor-constant-bufferize.mlir#L1))
 
     -   Bufferizes only `arith.constant` ops of `tensor` type.
     -   This is an example of setting up the legality so that only a subset of
-        `std.constant` ops get bufferized.
+        `arith.constant` ops get bufferized.
     -   This is an example of a pass that is not split along dialect
         subdivisions.
 
 ## How to write a finalizing bufferization pass
 
 The contract of a finalizing bufferization pass is that all tensors are gone
 from the program.
 

mlir/docs/DialectConversion.md

 struct MyTarget : public ConversionTarget {
   MyTarget(MLIRContext &ctx) : ConversionTarget(ctx) {
     //--------------------------------------------------------------------------
     // Marking an operation as Legal:
 
     /// Mark all operations within the LLVM dialect are legal.
     addLegalDialect<LLVMDialect>();
 
-    /// Mark `std.constant` op is always legal on this target.
-    addLegalOp<ConstantOp>();
+    /// Mark `arith.constant` op is always legal on this target.
+    addLegalOp<arith::ConstantOp>();
 
     //--------------------------------------------------------------------------
     // Marking an operation as dynamically legal.
 
     /// Mark all operations within Affine dialect have dynamic legality
     /// constraints.
     addDynamicallyLegalDialect<AffineDialect>([](Operation *op) { ... });
 

mlir/docs/Dialects/Affine.md

 
 Example:
 
 ```mlir
 For example, a DmaStartOp operation that transfers 256 elements of a memref
 '%src' in memory space 0 at indices [%i + 3, %j] to memref '%dst' in memory
 space 1 at indices [%k + 7, %l], would be specified as follows:
 
-  %num_elements = constant 256
+  %num_elements = arith.constant 256
   %idx = arith.constant 0 : index
   %tag = memref.alloc() : memref<1xi32, 4>
   affine.dma_start %src[%i + 3, %j], %dst[%k + 7, %l], %tag[%idx],
     %num_elements :
       memref<40x128xf32, 0>, memref<2x1024xf32, 1>, memref<1xi32, 2>
 
   If %stride and %num_elt_per_stride are specified, the DMA is expected to
   transfer %num_elt_per_stride elements every %stride elements apart from

mlir/docs/PassManagement.md

 *   `print-ir-before-all`
     *   Print the IR before every pass in the pipeline.
 
 ```shell
 $ mlir-opt foo.mlir -pass-pipeline='builtin.func(cse)' -print-ir-before=cse
 
 *** IR Dump Before CSE ***
 func @simple_constant() -> (i32, i32) {
-  %c1_i32 = constant 1 : i32
-  %c1_i32_0 = constant 1 : i32
+  %c1_i32 = arith.constant 1 : i32
+  %c1_i32_0 = arith.constant 1 : i32
   return %c1_i32, %c1_i32_0 : i32, i32
 }
 ```
 
 *   `print-ir-after=(comma-separated-pass-list)`
     *   Print the IR after each of the passes provided within the pass list.
 *   `print-ir-after-all`
     *   Print the IR after every pass in the pipeline.
 
 ```shell
 $ mlir-opt foo.mlir -pass-pipeline='builtin.func(cse)' -print-ir-after=cse
 
 *** IR Dump After CSE ***
 func @simple_constant() -> (i32, i32) {
-  %c1_i32 = constant 1 : i32
+  %c1_i32 = arith.constant 1 : i32
   return %c1_i32, %c1_i32 : i32, i32
 }
 ```
 
 *   `print-ir-after-change`
     *   Only print the IR after a pass if the pass mutated the IR. This helps to
         reduce the number of IR dumps for "uninteresting" passes.
     *   Note: Changes are detected by comparing a hash of the operation before
         and after the pass. This adds additional run-time to compute the hash of
         the IR, and in some rare cases may result in false-positives depending
         on the collision rate of the hash algorithm used.
     *   Note: This option should be used in unison with one of the other
         'print-ir-after' options above, as this option alone does not enable
         printing.
 
 ```shell
 $ mlir-opt foo.mlir -pass-pipeline='builtin.func(cse,cse)' -print-ir-after=cse -print-ir-after-change
 
 *** IR Dump After CSE ***
 func @simple_constant() -> (i32, i32) {
-  %c1_i32 = constant 1 : i32
+  %c1_i32 = arith.constant 1 : i32
   return %c1_i32, %c1_i32 : i32, i32
 }
 ```
 
 *   `print-ir-after-failure`
     *   Only print IR after a pass failure.
     *   This option should *not* be used with the other `print-ir-after` flags
         above.
 
 ```shell
 $ mlir-opt foo.mlir -pass-pipeline='builtin.func(cse,bad-pass)' -print-ir-failure
 
 *** IR Dump After BadPass Failed ***
 func @simple_constant() -> (i32, i32) {
-  %c1_i32 = constant 1 : i32
+  %c1_i32 = arith.constant 1 : i32
   return %c1_i32, %c1_i32 : i32, i32
 }
 ```
 
 *   `print-ir-module-scope`
     *   Always print the top-level module operation, regardless of pass type or
         operation nesting level.
     *   Note: Printing at module scope should only be used when multi-threading
         is disabled(`-mlir-disable-threading`)
 
 ```shell
 $ mlir-opt foo.mlir -mlir-disable-threading -pass-pipeline='builtin.func(cse)' -print-ir-after=cse -print-ir-module-scope
 
 *** IR Dump After CSE ***  ('builtin.func' operation: @bar)
 func @bar(%arg0: f32, %arg1: f32) -> f32 {
   ...
 }
 
 func @simple_constant() -> (i32, i32) {
-  %c1_i32 = constant 1 : i32
-  %c1_i32_0 = constant 1 : i32
+  %c1_i32 = arith.constant 1 : i32
+  %c1_i32_0 = arith.constant 1 : i32
   return %c1_i32, %c1_i32_0 : i32, i32
 }
 
 *** IR Dump After CSE ***  ('builtin.func' operation: @simple_constant)
 func @bar(%arg0: f32, %arg1: f32) -> f32 {
   ...
 }
 
 func @simple_constant() -> (i32, i32) {
-  %c1_i32 = constant 1 : i32
+  %c1_i32 = arith.constant 1 : i32
   return %c1_i32, %c1_i32 : i32, i32
 }
 ```
 
 ## Crash and Failure Reproduction
 
 The [pass manager](#pass-manager) in MLIR contains a builtin mechanism to
 generate reproducibles in the event of a crash, or a

mlir/docs/Passes.md

 ## `async` Dialect Passes
 
 [include "AsyncPasses.md"]
 
 ## `affine` Dialect Passes
 
 [include "AffinePasses.md"]
 
+## `arith` Dialect Passes
+
+[include "ArithmeticPasses.md"]
+
 ## `gpu` Dialect Passes
 
 [include "GPUPasses.md"]
 
 ## `linalg` Dialect Passes
 
 [include "LinalgPasses.md"]
 

mlir/docs/Rationale/MLIRForGraphAlgorithms.md

 
 ```mlir
   // RUN: mlir-opt %s -canonicalize | FileCheck %s
   func @test_subi_zero_cfg(%arg0: i32) -> i32 {
     %y = arith.subi %arg0, %arg0 : i32
     return %y: i32
   }
   // CHECK-LABEL: func @test_subi_zero_cfg(%arg0: i32)
-  // CHECK-NEXT: %c0_i32 = constant 0 : i32
+  // CHECK-NEXT: %c0_i32 = arith.constant 0 : i32
   // CHECK-NEXT: return %c0
 ```
 
 The "CHECK" comments are interpreted by the
 [LLVM FileCheck tool](https://llvm.org/docs/CommandGuide/FileCheck.html), which
 is sort of like a really advanced grep. This test is fully self-contained: it
 feeds the input into the [canonicalize pass](../Canonicalization.md), and checks
 that the output matches the CHECK lines. See the `test/Transforms` directory for

mlir/examples/standalone/include/Standalone/StandaloneOps.td

let description = [{
         operation in a dialect. It uses an operation trait to declare that it
         has no side effects.
 
         This operation takes an integer argument and returns an integer.
 
         Example:
 
         ```mlir
-        %0 = constant 2 : i32
+        %0 = arith.constant 2 : i32
         // Apply the foo operation to %0
         %1 = standalone.foo %0 : i32
         ```
     }];
 
     let arguments = (ins I32:$input);
     let results = (outs I32:$res);
 
     let assemblyFormat = [{
         $input attr-dict `:` type($input)
     }];
 }
 
 #endif // STANDALONE_OPS

mlir/include/mlir/Dialect/Affine/IR/AffineOps.h

 /// specified. The value of 'num_elements' must be a multiple of
 /// 'number_of_elements_per_stride'. If the source and destination locations
 /// overlap the behavior of this operation is not defined.
 //
 // For example, an AffineDmaStartOp operation that transfers 256 elements of a
 // memref '%src' in memory space 0 at indices [%i + 3, %j] to memref '%dst' in
 // memory space 1 at indices [%k + 7, %l], would be specified as follows:
 //
-//   %num_elements = constant 256
-//   %idx = constant 0 : index
+//   %num_elements = arith.constant 256
+//   %idx = arith.constant 0 : index
 //   %tag = alloc() : memref<1xi32, 4>
 //   affine.dma_start %src[%i + 3, %j], %dst[%k + 7, %l], %tag[%idx],
 //     %num_elements :
 //       memref<40x128xf32, 0>, memref<2x1024xf32, 1>, memref<1xi32, 2>
 //
 //   If %stride and %num_elt_per_stride are specified, the DMA is expected to
 //   transfer %num_elt_per_stride elements every %stride elements apart from
 //   memory space 0 until %num_elements are transferred.

mlir/include/mlir/Dialect/Affine/IR/AffineOps.td

let description = [{
     initial values of the loop-carried variables (corresponding to the SSA
     operands) will be the op's results.
 
     For example, to sum-reduce a memref:
 
      ```mlir
     func @reduce(%buffer: memref<1024xf32>) -> (f32) {
       // Initial sum set to 0.
-      %sum_0 = constant 0.0 : f32
+      %sum_0 = arith.constant 0.0 : f32
       // iter_args binds initial values to the loop's region arguments.
       %sum = affine.for %i = 0 to 10 step 2
           iter_args(%sum_iter = %sum_0) -> (f32) {
         %t = affine.load %buffer[%i] : memref<1024xf32>
         %sum_next = arith.addf %sum_iter, %t : f32
         // Yield current iteration sum to next iteration %sum_iter or to %sum
         // if final iteration.
         affine.yield %sum_next : f32
let description = [{
     #interior = affine_set<(i, j) : (i - 1 >= 0, j - 1 >= 0,  10 - i >= 0, 10 - j >= 0)> (%i, %j)
     func @pad_edges(%I : memref<10x10xf32>) -> (memref<12x12xf32) {
       %O = alloc memref<12x12xf32>
       affine.parallel (%i, %j) = (0, 0) to (12, 12) {
         %1 = affine.if #interior (%i, %j) {
           %2 = load %I[%i - 1, %j - 1] : memref<10x10xf32>
           affine.yield %2
         } else {
-          %2 = constant 0.0 : f32
+          %2 = arith.constant 0.0 : f32
           affine.yield %2 : f32
         }
         affine.store %1, %O[%i, %j] : memref<12x12xf32>
       }
       return %O
     }
     ```
   }];

mlir/include/mlir/Dialect/Affine/Passes.td

let description = [{
     This pass performs store to load forwarding and redundant load elimination
     for affine memref accesses and potentially eliminates the entire memref
     if all its accesses are forwarded.
 
     Input
 
     ```mlir
     func @store_load_affine_apply() -> memref<10x10xf32> {
-      %cf7 = constant 7.0 : f32
+      %cf7 = arith.constant 7.0 : f32
       %m = alloc() : memref<10x10xf32>
       affine.for %i0 = 0 to 10 {
         affine.for %i1 = 0 to 10 {
           affine.store %cf7, %m[%i0, %i1] : memref<10x10xf32>
           %v0 = affine.load %m[%i0, %i1] : memref<10x10xf32>
           %v1 = arith.addf %v0, %v0 : f32
         }
       }
       return %m : memref<10x10xf32>
     }
     ```
 
     Output
 
     ```mlir
     module {
       func @store_load_affine_apply() -> memref<10x10xf32> {
-        %cst = constant 7.000000e+00 : f32
+        %cst = arith.constant 7.000000e+00 : f32
         %0 = alloc() : memref<10x10xf32>
         affine.for %arg0 = 0 to 10 {
           affine.for %arg1 = 0 to 10 {
             affine.store %cst, %0[%arg0, %arg1] : memref<10x10xf32>
             %1 = arith.addf %cst, %cst : f32
           }
         }
         return %0 : memref<10x10xf32>

mlir/include/mlir/Dialect/Affine/Utils.h

 ///
 /// and VectorizationStrategy = 'vectorSizes = {8, 4}', 'loopToVectorDim =
 /// {{i1->0}, {i2->1}}', SuperVectorizer will generate:
 ///
 ///  func @vec2d(%arg0: memref<64x128x512xf32>, %arg1: memref<64x128x512xf32>) {
 ///    affine.for %arg2 = 0 to 64 {
 ///      affine.for %arg3 = 0 to 128 step 8 {
 ///        affine.for %arg4 = 0 to 512 step 4 {
-///          %cst = constant 0.000000e+00 : f32
+///          %cst = arith.constant 0.000000e+00 : f32
 ///          %0 = vector.transfer_read %arg0[%arg2, %arg3, %arg4], %cst : ...
 ///          vector.transfer_write %0, %arg1[%arg2, %arg3, %arg4] : ...
 ///        }
 ///      }
 ///    }
 ///    return
 ///  }
 // TODO: Hoist to a VectorizationStrategy.cpp when appropriate.

mlir/include/mlir/Dialect/MemRef/IR/MemRefOps.td

let description = [{
     If the dimension index is out of bounds the behavior is undefined.
 
     The specified memref type is that of the first operand.
 
     Example:
 
     ```mlir
     // Always returns 4, can be constant folded:
-    %c0 = constant 0 : index
+    %c0 = arith.constant 0 : index
     %x = memref.dim %A, %c0 : memref<4 x ? x f32>
 
     // Returns the dynamic dimension of %A.
-    %c1 = constant 1 : index
+    %c1 = arith.constant 1 : index
     %y = memref.dim %A, %c1 : memref<4 x ? x f32>
 
     // Equivalent generic form:
     %x = "memref.dim"(%A, %c0) : (memref<4 x ? x f32>, index) -> index
     %y = "memref.dim"(%A, %c1) : (memref<4 x ? x f32>, index) -> index
     ```
   }];
 
let description = [{
     source and destination locations overlap the behavior of this operation is
     not defined.
 
     For example, a DmaStartOp operation that transfers 256 elements of a memref
     '%src' in memory space 0 at indices [%i, %j] to memref '%dst' in memory
     space 1 at indices [%k, %l], would be specified as follows:
 
     ```mlir
-    %num_elements = constant 256
-    %idx = constant 0 : index
+    %num_elements = arith.constant 256
+    %idx = arith.constant 0 : index
     %tag = alloc() : memref<1 x i32, (d0) -> (d0), 4>
     dma_start %src[%i, %j], %dst[%k, %l], %num_elements, %tag[%idx] :
       memref<40 x 128 x f32>, (d0) -> (d0), 0>,
       memref<2 x 1024 x f32>, (d0) -> (d0), 1>,
       memref<1 x i32>, (d0) -> (d0), 2>
     ```
 
     If %stride and %num_elt_per_stride are specified, the DMA is expected to

mlir/include/mlir/Dialect/SCF/SCFOps.td

let description = [{
     }
 
     affine.for %i = 0 to 100 {
       "foo"() : () -> ()
       %v = scf.execute_region -> i64 {
         cond_br %cond, ^bb1, ^bb2
 
       ^bb1:
-        %c1 = constant 1 : i64
+        %c1 = arith.constant 1 : i64
         br ^bb3(%c1 : i64)
 
       ^bb2:
-        %c2 = constant 2 : i64
+        %c2 = arith.constant 2 : i64
         br ^bb3(%c2 : i64)
 
       ^bb3(%x : i64):
         scf.yield %x : i64
       }
       "bar"(%v) : (i64) -> ()
     }
     ```
let description = [{
 
     "scf.for" results hold the final values after the last iteration.
     For example, to sum-reduce a memref:
 
     ```mlir
     func @reduce(%buffer: memref<1024xf32>, %lb: index,
                  %ub: index, %step: index) -> (f32) {
       // Initial sum set to 0.
-      %sum_0 = constant 0.0 : f32
+      %sum_0 = arith.constant 0.0 : f32
       // iter_args binds initial values to the loop's region arguments.
       %sum = scf.for %iv = %lb to %ub step %step
           iter_args(%sum_iter = %sum_0) -> (f32) {
         %t = load %buffer[%iv] : memref<1024xf32>
         %sum_next = arith.addf %sum_iter, %t : f32
         // Yield current iteration sum to next iteration %sum_iter or to %sum
         // if final iteration.
         scf.yield %sum_next : f32
       }
       return %sum : f32
     }
     ```
 
     If the "scf.for" defines any values, a yield must be explicitly present.
     The number and types of the "scf.for" results must match the initial
     values in the "iter_args" binding and the yield operands.
 
     Another example with a nested "scf.if" (see "scf.if" for details) to
     perform conditional reduction:
 
     ```mlir
     func @conditional_reduce(%buffer: memref<1024xf32>, %lb: index,
                              %ub: index, %step: index) -> (f32) {
-      %sum_0 = constant 0.0 : f32
-      %c0 = constant 0.0 : f32
+      %sum_0 = arith.constant 0.0 : f32
+      %c0 = arith.constant 0.0 : f32
       %sum = scf.for %iv = %lb to %ub step %step
           iter_args(%sum_iter = %sum_0) -> (f32) {
         %t = load %buffer[%iv] : memref<1024xf32>
         %cond = arith.cmpf "ugt", %t, %c0 : f32
         %sum_next = scf.if %cond -> (f32) {
           %new_sum = arith.addf %sum_iter, %t : f32
           scf.yield %new_sum : f32
         } else {
let description = [{
 
     The body region must contain exactly one block that terminates with
     "scf.yield" without operands. Parsing ParallelOp will create such a region
     and insert the terminator when it is absent from the custom format.
 
     Example:
 
     ```mlir
-    %init = constant 0.0 : f32
+    %init = arith.constant 0.0 : f32
     scf.parallel (%iv) = (%lb) to (%ub) step (%step) init (%init) -> f32 {
       %elem_to_reduce = load %buffer[%iv] : memref<100xf32>
       scf.reduce(%elem_to_reduce) : f32 {
         ^bb0(%lhs : f32, %rhs: f32):
           %res = arith.addf %lhs, %rhs : f32
           scf.reduce.return %res : f32
       }
     }
let description = [{
     reduction may be non-deterministic unless the operation is associative and
     commutative.
 
     The result of the reduce operation's body must have the same type as the
     operands and associated result value of the parallel loop operation.
     Example:
 
     ```mlir
-    %operand = constant 1.0 : f32
+    %operand = arith.constant 1.0 : f32
     scf.reduce(%operand) : f32 {
       ^bb0(%lhs : f32, %rhs: f32):
         %res = arith.addf %lhs, %rhs : f32
         scf.reduce.return %res : f32
     }
     ```
   }];
 

mlir/include/mlir/Dialect/SparseTensor/Transforms/Passes.td

let description = [{
     relatively simple. In principle, these primitives could also be
     converted to actual elaborate IR code that implements the primitives
     on the selected sparse tensor storage schemes.
 
     Example of the conversion:
 
     ```mlir
       Before:
-        %c1 = constant 1 : index
+        %c1 = arith.constant 1 : index
         %0 = sparse_tensor.pointers %arg0, %c1
           : tensor<8x8xf32, #sparse_tensor.encoding<{
               dimLevelType = [ "dense", "compressed" ],
               pointerBitWidth = 0,
               indexBitWidth = 0
             }>> to memref<?xindex>
 
       After:
-        %c1 = constant 1 : index
+        %c1 = arith.constant 1 : index
         %0 = call @sparsePointers(%arg0, %c1) : (!llvm.ptr<i8>, index) -> memref<?xindex>
     ```
   }];
   let constructor = "mlir::createSparseTensorConversionPass()";
   let dependentDialects = [
     "arith::ArithmeticDialect",
     "LLVM::LLVMDialect",
     "memref::MemRefDialect",
     "scf::SCFDialect",
     "sparse_tensor::SparseTensorDialect",
     "vector::VectorDialect",
   ];
 }
 
 #endif // MLIR_DIALECT_SPARSETENSOR_TRANSFORMS_PASSES

mlir/include/mlir/Dialect/StandardOps/IR/Ops.td

let description = [{
     performed. No side-effecting ops are allowed in the body of
     `GenericAtomicRMWOp`.
 
     Example:
 
     ```mlir
     %x = generic_atomic_rmw %I[%i] : memref<10xf32> {
       ^bb0(%current_value : f32):
-        %c1 = constant 1.0 : f32
+        %c1 = arith.constant 1.0 : f32
         %inc = arith.addf %c1, %current_value : f32
         atomic_yield %inc : f32
     }
     ```
   }];
 
   let arguments = (ins
       MemRefOf<[AnySignlessInteger, AnyFloat]>:$memref,
let description = [{
     The `constant` operation produces an SSA value equal to some constant
     specified by an attribute. This is the way that MLIR uses to form simple
     integer and floating point constants, as well as more exotic things like
     references to functions and tensor/vector constants.
 
     Example:
 
     ```mlir
-    // Integer constant
-    %1 = constant 42 : i32
+    // Complex constant
+    %1 = constant [1.0 : f32, 1.0 : f32] : complex<f32>
 
     // Reference to function @myfn.
-    %3 = constant @myfn : (tensor<16xf32>, f32) -> tensor<16xf32>
+    %2 = constant @myfn : (tensor<16xf32>, f32) -> tensor<16xf32>
 
     // Equivalent generic forms
-    %1 = "std.constant"() {value = 42 : i32} : () -> i32
-    %3 = "std.constant"() {value = @myfn}
+    %1 = "std.constant"() {value = [1.0 : f32, 1.0 : f32] : complex<f32>}
+       : () -> complex<f32>
+    %2 = "std.constant"() {value = @myfn}
        : () -> ((tensor<16xf32>, f32) -> tensor<16xf32>)
     ```
 
     MLIR does not allow direct references to functions in SSA operands because
     the compiler is multithreaded, and disallowing SSA values to directly
     reference a function simplifies this
     ([rationale](../Rationale/Rationale.md#multithreading-the-compiler)).
   }];

mlir/include/mlir/Dialect/Tensor/IR/TensorOps.td

let description = [{
     If the dimension index is out of bounds, the behavior is undefined.
 
     The specified tensor type is that of the first operand.
 
     Example:
 
     ```mlir
     // Always returns 4, can be constant folded:
-    %c0 = constant 0 : index
+    %c0 = arith.constant 0 : index
     %x = tensor.dim %A, %c0 : tensor<4x?xf32>
 
     // Returns the dynamic dimension of %A.
-    %c1 = constant 1 : index
+    %c1 = arith.constant 1 : index
     %y = tensor.dim %A, %c1 : memref<4x?xf32>
 
     // Equivalent generic form:
     %x = "tensor.dim"(%A, %c0) : (memref<4x?xf32>, index) -> index
     %y = "tensor.dim"(%A, %c1) : (memref<4x?xf32>, index) -> index
     ```
   }];
 

mlir/include/mlir/Dialect/Vector/VectorOps.td

let description = [{
     The source operand is duplicated over all the missing leading dimensions
     and stretched over the trailing dimensions where the source has a non-equal
     dimension of 1. These rules imply that any scalar broadcast (k=0) to any
     shaped vector with the same element type is always legal.
 
     Example:
 
     ```mlir
-    %0 = constant 0.0 : f32
+    %0 = arith.constant 0.0 : f32
     %1 = vector.broadcast %0 : f32 to vector<16xf32>
     %2 = vector.broadcast %1 : vector<16xf32> to vector<4x16xf32>
     ```
   }];
   let extraClassDeclaration = [{
     Type getSourceType() { return source().getType(); }
     VectorType getVectorType() {
       return vector().getType().cast<VectorType>();
let description = [{
     scalar at that position. Note that this instruction resembles
     vector.extract, but is restricted to 1-D vectors and relaxed
     to dynamic indices. It is meant to be closer to LLVM's version:
     https://llvm.org/docs/LangRef.html#extractelement-instruction
 
     Example:
 
     ```mlir
-    %c = constant 15 : i32
+    %c = arith.constant 15 : i32
     %1 = vector.extractelement %0[%c : i32]: vector<16xf32>
     ```
   }];
   let assemblyFormat = [{
     $vector `[` $position `:` type($position) `]` attr-dict `:` type($vector)
   }];
 
   let builders = [
let description = [{
     position.  Note that this instruction resembles vector.insert, but
     is restricted to 1-D vectors and relaxed to dynamic indices. It is
     meant to be closer to LLVM's version:
     https://llvm.org/docs/LangRef.html#insertelement-instruction
 
     Example:
 
     ```mlir
-    %c = constant 15 : i32
-    %f = constant 0.0f : f32
+    %c = arith.constant 15 : i32
+    %f = arith.constant 0.0f : f32
     %1 = vector.insertelement %f, %0[%c : i32]: vector<16xf32>
     ```
   }];
   let assemblyFormat = [{
     $source `,` $dest `[` $position `:` type($position) `]` attr-dict `:`
     type($result)
   }];
 
let description = [{
     %view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>>
     for %i = 0 to 3 {
       affine.for %j = 0 to 4 {
         affine.for %k = 0 to 5 {
           %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] :
             memref<?x?x?x?xf32>
           store %tmp[%i, %j, %k] : vector<3x4x5xf32>
     }}}
-    %c0 = constant 0 : index
+    %c0 = arith.constant 0 : index
     %vec = load %view_in_tmp[%c0] : vector<3x4x5xf32>
     ```
 
     On a GPU one could then map `i`, `j`, `k` to blocks and threads. Notice that
     the temporary storage footprint is `3 * 5` values but `3 * 4 * 5` values are
     actually transferred between `%A` and `%tmp`.
 
     Alternatively, if a notional vector broadcast operation were available, the
     lowered code would resemble:
 
     ```mlir
     // %expr1, %expr2, %expr3, %expr4 defined before this point
     %tmp = alloc() : vector<3x4x5xf32>
     %view_in_tmp = "element_type_cast"(%tmp) : memref<1xvector<3x4x5xf32>>
     for %i = 0 to 3 {
       affine.for %k = 0 to 5 {
         %a = load %A[%expr1 + %k, %expr2, %expr3 + %i, %expr4] :
           memref<?x?x?x?xf32>
         store %tmp[%i, 0, %k] : vector<3x4x5xf32>
     }}
-    %c0 = constant 0 : index
+    %c0 = arith.constant 0 : index
     %tmpvec = load %view_in_tmp[%c0] : vector<3x4x5xf32>
     %vec = broadcast %tmpvec, 1 : vector<3x4x5xf32>
     ```
 
     where `broadcast` broadcasts from element 0 to all others along the
     specified dimension. This time, the temporary storage footprint is `3 * 5`
     values which is the same amount of data as the `3 * 5` values transferred.
     An additional `1` broadcast is required. On a GPU this broadcast could be
     implemented using a warp-shuffle if loop `j` were mapped to `threadIdx.x`.
 
     Syntax
     ```
     operation ::= ssa-id `=` `vector.transfer_read` ssa-use-list
       `{` attribute-entry `} :` memref-type `,` vector-type
     ```
 
     Example:
 
     ```mlir
     // Read the slice `%A[%i0, %i1:%i1+256, %i2:%i2+32]` into vector<32x256xf32>
     // and pad with %f0 to handle the boundary case:
-    %f0 = constant 0.0f : f32
+    %f0 = arith.constant 0.0f : f32
     for %i0 = 0 to %0 {
       affine.for %i1 = 0 to %1 step 256 {
         affine.for %i2 = 0 to %2 step 32 {
           %v = vector.transfer_read %A[%i0, %i1, %i2], (%f0)
                {permutation_map: (d0, d1, d2) -> (d2, d1)} :
                memref<?x?x?xf32>, vector<32x256xf32>
     }}}
 
def Vector_PrintOp :
   let summary = "print operation (for testing and debugging)";
   let description = [{
     Prints the source vector (or scalar) to stdout in human readable
     format (for testing and debugging). No return value.
 
     Example:
 
     ```mlir
-    %0 = constant 0.0 : f32
+    %0 = arith.constant 0.0 : f32
     %1 = vector.broadcast %0 : f32 to vector<4xf32>
     vector.print %1 : vector<4xf32>
 
     when lowered to LLVM, the vector print is unrolled into
     elementary printing method calls that at runtime will yield
 
     ( 0.0, 0.0, 0.0, 0.0 )
 

mlir/include/mlir/Dialect/Vector/VectorTransforms.h

private:
   /// Options to control the vector patterns.
   vector::VectorTransformsOptions vectorTransformOptions;
   FilterConstraintType filter;
 };
 
 /// Progressive lowering of a `vector.contract %a, %b, %c` with row-major matmul
 /// semantics to an output-size-unrolled sequence:
 /// ```
-///    %out = constant ... : vector<MxNxelt_type>
+///    %out = arith.constant ... : vector<MxNxelt_type>
 ///    %bt = vector.transpose %b, [1, 0]
 ///    %aRow0 = vector.extract %a[0]
 ///    %btRow0 = vector.extract %bt[0]
 ///    %c00 = vector.reduce %atRow0, %bRow0
 ///    %out00 = vector.insert %c00, %out[0, 0]
 ///    ...
 ///    %aRowLast = vector.extract %at[M-1]
 ///    %btRowLast = vector.extract %b[N-1]

mlir/include/mlir/Dialect/Vector/VectorUtils.h

 ///
 /// Meaning that vector.transfer_read will be responsible for reading the slice:
 /// `%arg0[%i4, %i5:%15+256, %i3:%i3+32]` into vector<32x256xf32>.
 ///
 /// Example 2:
 /// The following MLIR snippet:
 ///
 /// ```mlir
-///    %cst0 = constant 0 : index
+///    %cst0 = arith.constant 0 : index
 ///    affine.for %i0 = 0 to %0 {
 ///      %a0 = load %arg0[%cst0, %cst0] : memref<?x?xf32>
 ///    }
 /// ```
 ///
 /// may vectorize with {permutation_map: (d0) -> (0)} into:
 ///
 /// ```mlir

mlir/include/mlir/Transforms/Passes.td

let description = [{
     which a source slice should be materialized in the designation slice.
 
     Example 1: Producer-consumer fusion.
     Input:
     ```mlir
     func @producer_consumer_fusion(%arg0: memref<10xf32>, %arg1: memref<10xf32>) {
       %0 = alloc() : memref<10xf32>
       %1 = alloc() : memref<10xf32>
-      %cst = constant 0.000000e+00 : f32
+      %cst = arith.constant 0.000000e+00 : f32
       affine.for %arg2 = 0 to 10 {
         affine.store %cst, %0[%arg2] : memref<10xf32>
         affine.store %cst, %1[%arg2] : memref<10xf32>
       }
       affine.for %arg2 = 0 to 10 {
         %2 = affine.load %0[%arg2] : memref<10xf32>
         %3 = arith.addf %2, %2 : f32
         affine.store %3, %arg0[%arg2] : memref<10xf32>
       }
       affine.for %arg2 = 0 to 10 {
         %2 = affine.load %1[%arg2] : memref<10xf32>
         %3 = arith.mulf %2, %2 : f32
         affine.store %3, %arg1[%arg2] : memref<10xf32>
       }
       return
     }
     ```
     Output:
     ```mlir
     func @producer_consumer_fusion(%arg0: memref<10xf32>, %arg1: memref<10xf32>) {
       %0 = alloc() : memref<1xf32>
       %1 = alloc() : memref<1xf32>
-      %cst = constant 0.000000e+00 : f32
+      %cst = arith.constant 0.000000e+00 : f32
       affine.for %arg2 = 0 to 10 {
         affine.store %cst, %0[0] : memref<1xf32>
         affine.store %cst, %1[0] : memref<1xf32>
         %2 = affine.load %1[0] : memref<1xf32>
         %3 = arith.mulf %2, %2 : f32
         affine.store %3, %arg1[%arg2] : memref<10xf32>
         %4 = affine.load %0[0] : memref<1xf32>
         %5 = arith.addf %4, %4 : f32
let description = [{
 
     Input
 
     ```mlir
     func @pipelinedatatransfer() {
       %0 = alloc() : memref<256xf32>
       %1 = alloc() : memref<32xf32, 1>
       %2 = alloc() : memref<1xf32>
-      %c0 = constant 0 : index
-      %c128 = constant 128 : index
+      %c0 = arith.constant 0 : index
+      %c128 = arith.constant 128 : index
       affine.for %i0 = 0 to 8 {
         affine.dma_start %0[%i0], %1[%i0], %2[%c0], %c128 : memref<256xf32>, memref<32xf32, 1>, memref<1xf32>
         affine.dma_wait %2[%c0], %c128 : memref<1xf32>
         %3 = affine.load %1[%i0] : memref<32xf32, 1>
         %4 = "compute"(%3) : (f32) -> f32
         affine.store %4, %1[%i0] : memref<32xf32, 1>
       }
       return
     }
     ```
 
     Output
 
     ```mlir
     module {
       func @pipelinedatatransfer() {
-        %c8 = constant 8 : index
-        %c0 = constant 0 : index
+        %c8 = arith.constant 8 : index
+        %c0 = arith.constant 0 : index
         %0 = alloc() : memref<256xf32>
-        %c0_0 = constant 0 : index
-        %c128 = constant 128 : index
+        %c0_0 = arith.constant 0 : index
+        %c128 = arith.constant 128 : index
         %1 = alloc() : memref<2x32xf32, 1>
         %2 = alloc() : memref<2x1xf32>
         affine.dma_start %0[%c0], %1[%c0 mod 2, %c0], %2[%c0 mod 2, symbol(%c0_0)], %c128 : memref<256xf32>, memref<2x32xf32, 1>, memref<2x1xf32>
         affine.for %arg0 = 1 to 8 {
           affine.dma_start %0[%arg0], %1[%arg0 mod 2, %arg0], %2[%arg0 mod 2, symbol(%c0_0)], %c128 : memref<256xf32>, memref<2x32xf32, 1>, memref<2x1xf32>
           %8 = affine.apply #map3(%arg0)
           %9 = affine.apply #map4(%8)
           %10 = affine.apply #map4(%8)
let description = [{
 
     Input
 
     ```
     #linear8 = affine_map<(i, j) -> (i * 8 + j)>
     func @linearize(%arg0: memref<8x8xi32, #linear8>,
                     %arg1: memref<8x8xi32, #linear8>,
                     %arg2: memref<8x8xi32, #linear8>) {
-      %c8 = constant 8 : index
-      %c0 = constant 0 : index
-      %c1 = constant 1 : index
+      %c8 = arith.constant 8 : index
+      %c0 = arith.constant 0 : index
+      %c1 = arith.constant 1 : index
       affine.for %arg3 = %c0 to %c8  {
       affine.for %arg4 = %c0 to %c8  {
         affine.for %arg5 = %c0 to %c8 {
           %0 = affine.load %arg0[%arg3, %arg5] : memref<8x8xi32, #linear8>
           %1 = affine.load %arg1[%arg5, %arg4] : memref<8x8xi32, #linear8>
           %2 = affine.load %arg2[%arg3, %arg4] : memref<8x8xi32, #linear8>
           %3 = arith.muli %0, %1 : i32
           %4 = arith.addi %2, %3 : i32
           affine.store %4, %arg2[%arg3, %arg4] : memref<8x8xi32, #linear8>
         }
       }
       }
       return
     }
     ```
 
     Output
 
     ```mlir
     func @linearize(%arg0: memref<64xi32>,
                     %arg1: memref<64xi32>,
                     %arg2: memref<64xi32>) {
-    %c8 = constant 8 : index
-    %c0 = constant 0 : index
+    %c8 = arith.constant 8 : index
+    %c0 = arith.constant 0 : index
     affine.for %arg3 = %c0 to %c8 {
       affine.for %arg4 = %c0 to %c8 {
         affine.for %arg5 = %c0 to %c8 {
           %0 = affine.load %arg0[%arg3 * 8 + %arg5] : memref<64xi32>
           %1 = affine.load %arg1[%arg5 * 8 + %arg4] : memref<64xi32>
           %2 = affine.load %arg2[%arg3 * 8 + %arg4] : memref<64xi32>
           %3 = arith.muli %0, %1 : i32
           %4 = arith.addi %2, %3 : i32

mlir/lib/Conversion/ArithmeticToSPIRV/ArithmeticToSPIRV.cpp

 
 LogicalResult ConstantCompositeOpPattern::matchAndRewrite(
     arith::ConstantOp constOp, OpAdaptor adaptor,
     ConversionPatternRewriter &rewriter) const {
   auto srcType = constOp.getType().dyn_cast<ShapedType>();
   if (!srcType)
     return failure();
 
-  // std.constant should only have vector or tenor types.
+  // arith.constant should only have vector or tenor types.
   assert((srcType.isa<VectorType, RankedTensorType>()));
 
   auto dstType = getTypeConverter()->convertType(srcType);
   if (!dstType)
     return failure();
 
   auto dstElementsAttr = constOp.value().dyn_cast<DenseElementsAttr>();
   ShapedType dstAttrType = dstElementsAttr.getType();
if (srcType.isa<FloatType>()) {
     }
 
     rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constOp, dstType, dstAttr);
     return success();
   }
 
   // Bool type.
   if (srcType.isInteger(1)) {
-    // std.constant can use 0/1 instead of true/false for i1 values. We need to
-    // handle that here.
+    // arith.constant can use 0/1 instead of true/false for i1 values. We need
+    // to handle that here.
     auto dstAttr = convertBoolAttr(constOp.value(), rewriter);
     if (!dstAttr)
       return failure();
     rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constOp, dstType, dstAttr);
     return success();
   }
 
   // IndexType or IntegerType. Index values are converted to 32-bit integer

mlir/lib/Conversion/ShapeToStandard/ShapeToStandard.cpp

 /// compares their size and, if equal, checks every extent for equality.
 ///
 /// Example:
 ///
 /// %result = shape.shape_eq %a, %b : tensor<?xindex>, tensor<?xindex>
 ///
 /// becomes
 ///
-/// %c0 = constant 0 : index
+/// %c0 = arith.constant 0 : index
 /// %0 = dim %arg0, %c0 : tensor<?xindex>
 /// %1 = dim %arg1, %c0 : tensor<?xindex>
 /// %2 = arith.cmpi "eq", %0, %1 : index
 /// %result = scf.if %2 -> (i1) {
 ///   %c1 = arith.constant 1 : index
 ///   %true = arith.constant true
 ///   %4 = scf.for %arg2 = %c0 to %0 step %c1 iter_args(%arg3 = %true) -> (i1) {
 ///     %5 = tensor.extract %arg0[%arg2] : tensor<?xindex>

mlir/lib/Conversion/StandardToSPIRV/StandardToSPIRV.cpp

 #include "llvm/ADT/SetVector.h"
 #include "llvm/Support/Debug.h"
 
 #define DEBUG_TYPE "std-to-spirv-pattern"
 
 using namespace mlir;
 
 //===----------------------------------------------------------------------===//
-// Utility functions
-//===----------------------------------------------------------------------===//
-
-/// Converts the given `srcAttr` into a boolean attribute if it holds an
-/// integral value. Returns null attribute if conversion fails.
-static BoolAttr convertBoolAttr(Attribute srcAttr, Builder builder) {
-  if (auto boolAttr = srcAttr.dyn_cast<BoolAttr>())
-    return boolAttr;
-  if (auto intAttr = srcAttr.dyn_cast<IntegerAttr>())
-    return builder.getBoolAttr(intAttr.getValue().getBoolValue());
-  return BoolAttr();
-}
-
-/// Converts the given `srcAttr` to a new attribute of the given `dstType`.
-/// Returns null attribute if conversion fails.
-static IntegerAttr convertIntegerAttr(IntegerAttr srcAttr, IntegerType dstType,
-                                      Builder builder) {
-  // If the source number uses less active bits than the target bitwidth, then
-  // it should be safe to convert.
-  if (srcAttr.getValue().isIntN(dstType.getWidth()))
-    return builder.getIntegerAttr(dstType, srcAttr.getInt());
-
-  // XXX: Try again by interpreting the source number as a signed value.
-  // Although integers in the standard dialect are signless, they can represent
-  // a signed number. It's the operation decides how to interpret. This is
-  // dangerous, but it seems there is no good way of handling this if we still
-  // want to change the bitwidth. Emit a message at least.
-  if (srcAttr.getValue().isSignedIntN(dstType.getWidth())) {
-    auto dstAttr = builder.getIntegerAttr(dstType, srcAttr.getInt());
-    LLVM_DEBUG(llvm::dbgs() << "attribute '" << srcAttr << "' converted to '"
-                            << dstAttr << "' for type '" << dstType << "'\n");
-    return dstAttr;
-  }
-
-  LLVM_DEBUG(llvm::dbgs() << "attribute '" << srcAttr
-                          << "' illegal: cannot fit into target type '"
-                          << dstType << "'\n");
-  return IntegerAttr();
-}
-
-/// Converts the given `srcAttr` to a new attribute of the given `dstType`.
-/// Returns null attribute if `dstType` is not 32-bit or conversion fails.
-static FloatAttr convertFloatAttr(FloatAttr srcAttr, FloatType dstType,
-                                  Builder builder) {
-  // Only support converting to float for now.
-  if (!dstType.isF32())
-    return FloatAttr();
-
-  // Try to convert the source floating-point number to single precision.
-  APFloat dstVal = srcAttr.getValue();
-  bool losesInfo = false;
-  APFloat::opStatus status =
-      dstVal.convert(APFloat::IEEEsingle(), APFloat::rmTowardZero, &losesInfo);
-  if (status != APFloat::opOK || losesInfo) {
-    LLVM_DEBUG(llvm::dbgs()
-               << srcAttr << " illegal: cannot fit into converted type '"
-               << dstType << "'\n");
-    return FloatAttr();
-  }
-
-  return builder.getF32FloatAttr(dstVal.convertToFloat());
-}
-
-//===----------------------------------------------------------------------===//
 // Operation conversion
 //===----------------------------------------------------------------------===//
 
 // Note that DRR cannot be used for the patterns in this file: we may need to
 // convert type along the way, which requires ConversionPattern. DRR generates
 // normal RewritePattern.
 
 namespace {
 
-/// Converts composite std.constant operation to spv.Constant.
-class ConstantCompositeOpPattern final
-    : public OpConversionPattern<ConstantOp> {
-public:
-  using OpConversionPattern<ConstantOp>::OpConversionPattern;
-
-  LogicalResult
-  matchAndRewrite(ConstantOp constOp, OpAdaptor adaptor,
-                  ConversionPatternRewriter &rewriter) const override;
-};
-
-/// Converts scalar std.constant operation to spv.Constant.
-class ConstantScalarOpPattern final : public OpConversionPattern<ConstantOp> {
-public:
-  using OpConversionPattern<ConstantOp>::OpConversionPattern;
-
-  LogicalResult
-  matchAndRewrite(ConstantOp constOp, OpAdaptor adaptor,
-                  ConversionPatternRewriter &rewriter) const override;
-};
-
 /// Converts std.return to spv.Return.
 class ReturnOpPattern final : public OpConversionPattern<ReturnOp> {
 public:
   using OpConversionPattern<ReturnOp>::OpConversionPattern;
 
   LogicalResult
   matchAndRewrite(ReturnOp returnOp, OpAdaptor adaptor,
                   ConversionPatternRewriter &rewriter) const override;

 
 private:
   int64_t byteCountThreshold;
 };
 
 } // namespace
 
 //===----------------------------------------------------------------------===//
-// ConstantOp with composite type.
-//===----------------------------------------------------------------------===//
-
-// TODO: This probably should be split into the vector case and tensor case,
-// so that the tensor case can be moved to TensorToSPIRV conversion. But,
-// std.constant is for the standard dialect though.
-LogicalResult ConstantCompositeOpPattern::matchAndRewrite(
-    ConstantOp constOp, OpAdaptor adaptor,
-    ConversionPatternRewriter &rewriter) const {
-  auto srcType = constOp.getType().dyn_cast<ShapedType>();
-  if (!srcType)
-    return failure();
-
-  // std.constant should only have vector or tenor types.
-  assert((srcType.isa<VectorType, RankedTensorType>()));
-
-  auto dstType = getTypeConverter()->convertType(srcType);
-  if (!dstType)
-    return failure();
-
-  auto dstElementsAttr = constOp.value().dyn_cast<DenseElementsAttr>();
-  ShapedType dstAttrType = dstElementsAttr.getType();
-  if (!dstElementsAttr)
-    return failure();
-
-  // If the composite type has more than one dimensions, perform linearization.
-  if (srcType.getRank() > 1) {
-    if (srcType.isa<RankedTensorType>()) {
-      dstAttrType = RankedTensorType::get(srcType.getNumElements(),
-                                          srcType.getElementType());
-      dstElementsAttr = dstElementsAttr.reshape(dstAttrType);
-    } else {
-      // TODO: add support for large vectors.
-      return failure();
-    }
-  }
-
-  Type srcElemType = srcType.getElementType();
-  Type dstElemType;
-  // Tensor types are converted to SPIR-V array types; vector types are
-  // converted to SPIR-V vector/array types.
-  if (auto arrayType = dstType.dyn_cast<spirv::ArrayType>())
-    dstElemType = arrayType.getElementType();
-  else
-    dstElemType = dstType.cast<VectorType>().getElementType();
-
-  // If the source and destination element types are different, perform
-  // attribute conversion.
-  if (srcElemType != dstElemType) {
-    SmallVector<Attribute, 8> elements;
-    if (srcElemType.isa<FloatType>()) {
-      for (FloatAttr srcAttr : dstElementsAttr.getValues<FloatAttr>()) {
-        FloatAttr dstAttr =
-            convertFloatAttr(srcAttr, dstElemType.cast<FloatType>(), rewriter);
-        if (!dstAttr)
-          return failure();
-        elements.push_back(dstAttr);
-      }
-    } else if (srcElemType.isInteger(1)) {
-      return failure();
-    } else {
-      for (IntegerAttr srcAttr : dstElementsAttr.getValues<IntegerAttr>()) {
-        IntegerAttr dstAttr = convertIntegerAttr(
-            srcAttr, dstElemType.cast<IntegerType>(), rewriter);
-        if (!dstAttr)
-          return failure();
-        elements.push_back(dstAttr);
-      }
-    }
-
-    // Unfortunately, we cannot use dialect-specific types for element
-    // attributes; element attributes only works with builtin types. So we need
-    // to prepare another converted builtin types for the destination elements
-    // attribute.
-    if (dstAttrType.isa<RankedTensorType>())
-      dstAttrType = RankedTensorType::get(dstAttrType.getShape(), dstElemType);
-    else
-      dstAttrType = VectorType::get(dstAttrType.getShape(), dstElemType);
-
-    dstElementsAttr = DenseElementsAttr::get(dstAttrType, elements);
-  }
-
-  rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constOp, dstType,
-                                                 dstElementsAttr);
-  return success();
-}
-
-//===----------------------------------------------------------------------===//
-// ConstantOp with scalar type.
-//===----------------------------------------------------------------------===//
-
-LogicalResult ConstantScalarOpPattern::matchAndRewrite(
-    ConstantOp constOp, OpAdaptor adaptor,
-    ConversionPatternRewriter &rewriter) const {
-  Type srcType = constOp.getType();
-  if (!srcType.isIntOrIndexOrFloat())
-    return failure();
-
-  Type dstType = getTypeConverter()->convertType(srcType);
-  if (!dstType)
-    return failure();
-
-  // Floating-point types.
-  if (srcType.isa<FloatType>()) {
-    auto srcAttr = constOp.value().cast<FloatAttr>();
-    auto dstAttr = srcAttr;
-
-    // Floating-point types not supported in the target environment are all
-    // converted to float type.
-    if (srcType != dstType) {
-      dstAttr = convertFloatAttr(srcAttr, dstType.cast<FloatType>(), rewriter);
-      if (!dstAttr)
-        return failure();
-    }
-
-    rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constOp, dstType, dstAttr);
-    return success();
-  }
-
-  // Bool type.
-  if (srcType.isInteger(1)) {
-    // std.constant can use 0/1 instead of true/false for i1 values. We need to
-    // handle that here.
-    auto dstAttr = convertBoolAttr(constOp.value(), rewriter);
-    if (!dstAttr)
-      return failure();
-    rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constOp, dstType, dstAttr);
-    return success();
-  }
-
-  // IndexType or IntegerType. Index values are converted to 32-bit integer
-  // values when converting to SPIR-V.
-  auto srcAttr = constOp.value().cast<IntegerAttr>();
-  auto dstAttr =
-      convertIntegerAttr(srcAttr, dstType.cast<IntegerType>(), rewriter);
-  if (!dstAttr)
-    return failure();
-  rewriter.replaceOpWithNewOp<spirv::ConstantOp>(constOp, dstType, dstAttr);
-  return success();
-}
-
-//===----------------------------------------------------------------------===//
 // ReturnOp
 //===----------------------------------------------------------------------===//
 
 LogicalResult
 ReturnOpPattern::matchAndRewrite(ReturnOp returnOp, OpAdaptor adaptor,
                                  ConversionPatternRewriter &rewriter) const {
   if (returnOp.getNumOperands() > 1)
     return failure();
patterns.add<
       // Unary and binary patterns
       spirv::UnaryAndBinaryOpPattern<MaxFOp, spirv::GLSLFMaxOp>,
       spirv::UnaryAndBinaryOpPattern<MaxSIOp, spirv::GLSLSMaxOp>,
       spirv::UnaryAndBinaryOpPattern<MaxUIOp, spirv::GLSLUMaxOp>,
       spirv::UnaryAndBinaryOpPattern<MinFOp, spirv::GLSLFMinOp>,
       spirv::UnaryAndBinaryOpPattern<MinSIOp, spirv::GLSLSMinOp>,
       spirv::UnaryAndBinaryOpPattern<MinUIOp, spirv::GLSLUMinOp>,
 
-      // Constant patterns
-      ConstantCompositeOpPattern, ConstantScalarOpPattern,
-
       ReturnOpPattern, SelectOpPattern, SplatPattern>(typeConverter, context);
 }
 
 void populateTensorToSPIRVPatterns(SPIRVTypeConverter &typeConverter,
                                    int64_t byteCountThreshold,
                                    RewritePatternSet &patterns) {
   patterns.add<TensorExtractPattern>(typeConverter, patterns.getContext(),
                                      byteCountThreshold);
 }
 
 } // namespace mlir

mlir/lib/Dialect/Async/Transforms/AsyncToAsyncRuntime.cpp

 namespace {
 struct CoroMachinery {
   FuncOp func;
 
   // Async execute region returns a completion token, and an async value for
   // each yielded value.
   //
   //   %token, %result = async.execute -> !async.value<T> {
-  //     %0 = constant ... : T
+  //     %0 = arith.constant ... : T
   //     async.yield %0 : T
   //   }
   Value asyncToken; // token representing completion of the async region
   llvm::SmallVector<Value, 4> returnValues; // returned async values
 
   Value coroHandle; // coroutine handle (!async.coro.handle value)
   Block *entry;     // coroutine entry block
   Block *setError;  // switch completion token and all values to error state

mlir/lib/Dialect/Linalg/IR/LinalgOps.cpp

 }
 
 namespace {
 /// Change the type of the result of a `linalg.init_tensor` by making the result
 /// type statically sized along dimension that in the original operation where
 /// defined as dynamic, but the size was defined using a `constant` op. For
 /// example
 ///
-///  %c5 = constant 5: index
+///  %c5 = arith.constant 5: index
 ///  %0 = linalg.init_tensor [%arg0, %c5] : tensor<?x?xf32>
 ///
 ///  to
 ///
 ///  %0 = linalg.init_tensor [%arg0, 5] : tensor<?x5xf32>
 struct ReplaceStaticShapeDims : OpRewritePattern<InitTensorOp> {
   using OpRewritePattern<InitTensorOp>::OpRewritePattern;
 

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

 //     %j = linalg.index 1 : index
 //     <some operations that use %i, %j>
 // }: memref<50x100xf32>, memref<50x100xf32>
 //
 // After tiling pass with tiles sizes 10 and 25:
 //
 // #strided = (i, j)[s0, s1, s2] -> (i * s1 + s0 + j * s2)
 //
-// %c1 = constant 1 : index
-// %c0 = constant 0 : index
-// %c25 = constant 25 : index
-// %c10 = constant 10 : index
+// %c1 = arith.constant 1 : index
+// %c0 = arith.constant 0 : index
+// %c25 = arith.constant 25 : index
+// %c10 = arith.constant 10 : index
 // operand_dim_0 = dim %operand, 0 : memref<50x100xf32>
 // operand_dim_1 = dim %operand, 1 : memref<50x100xf32>
 // scf.for %k = %c0 to operand_dim_0 step %c10 {
 //   scf.for %l = %c0 to operand_dim_1 step %c25 {
 //     %4 = std.subview %operand[%k, %l][%c10, %c25][%c1, %c1]
 //       : memref<50x100xf32> to memref<?x?xf32, #strided>
 //     %5 = std.subview %result[%k, %l][%c10, %c25][%c1, %c1]
 //       : memref<50x100xf32> to memref<?x?xf32, #strided>

mlir/lib/Transforms/Utils/Utils.cpp

 /// - New map in the above example
 ///   #map0 = affine_map<(d0, d1) -> (d0)>
 ///   #map1 = affine_map<(d0, d1) -> (d1 ceildiv 32)>
 ///   #map2 = affine_map<(d0, d1) -> (32)>
 ///
 /// 4. Create AffineApplyOp to apply the new maps. The output of AffineApplyOp
 /// is used in dynamicSizes of new AllocOp.
 ///   %0 = dim %arg0, %c1 : memref<4x?xf32>
-///   %c4 = constant 4 : index
+///   %c4 = arith.constant 4 : index
 ///   %1 = affine.apply #map1(%c4, %0)
 ///   %2 = affine.apply #map2(%c4, %0)
 static void createNewDynamicSizes(MemRefType oldMemRefType,
                                   MemRefType newMemRefType, AffineMap map,
                                   memref::AllocOp *allocOp, OpBuilder b,
                                   SmallVectorImpl<Value> &newDynamicSizes) {
   // Create new input for AffineApplyOp.
   SmallVector<Value, 4> inAffineApply;

mlir/test/CAPI/ir.c

int testTypeID(MlirContext ctx) {
   MlirStringRef valueStringRef = mlirStringRefCreateFromCString("value");
 
   // Create a registered operation, which should have a type id.
   MlirAttribute indexZeroLiteral =
       mlirAttributeParseGet(ctx, mlirStringRefCreateFromCString("0 : index"));
   MlirNamedAttribute indexZeroValueAttr = mlirNamedAttributeGet(
       mlirIdentifierGet(ctx, valueStringRef), indexZeroLiteral);
   MlirOperationState constZeroState = mlirOperationStateGet(
-      mlirStringRefCreateFromCString("std.constant"), loc);
+      mlirStringRefCreateFromCString("arith.constant"), loc);
   mlirOperationStateAddResults(&constZeroState, 1, &indexType);
   mlirOperationStateAddAttributes(&constZeroState, 1, &indexZeroValueAttr);
   MlirOperation constZero = mlirOperationCreate(&constZeroState);
 
+  if (!mlirOperationVerify(constZero)) {
+    fprintf(stderr, "ERROR: Expected operation to verify correctly\n");
+    return 5;
+  }
+
   if (mlirOperationIsNull(constZero)) {
     fprintf(stderr, "ERROR: Expected registered operation to be present\n");
-    return 5;
+    return 6;
   }
 
   MlirTypeID registeredOpID = mlirOperationGetTypeID(constZero);
 
   if (mlirTypeIDIsNull(registeredOpID)) {
     fprintf(stderr,
             "ERROR: Expected registered operation type id to be present\n");
-    return 6;
+    return 7;
   }
 
   // Create an unregistered operation, which should not have a type id.
   mlirContextSetAllowUnregisteredDialects(ctx, true);
   MlirOperationState opState =
       mlirOperationStateGet(mlirStringRefCreateFromCString("dummy.op"), loc);
   MlirOperation unregisteredOp = mlirOperationCreate(&opState);
   if (mlirOperationIsNull(unregisteredOp)) {
     fprintf(stderr, "ERROR: Expected unregistered operation to be present\n");
-    return 7;
+    return 8;
   }
 
   MlirTypeID unregisteredOpID = mlirOperationGetTypeID(unregisteredOp);
 
   if (!mlirTypeIDIsNull(unregisteredOpID)) {
     fprintf(stderr,
             "ERROR: Expected unregistered operation type id to be null\n");
-    return 8;
+    return 9;
   }
 
   mlirOperationDestroy(constZero);
   mlirOperationDestroy(unregisteredOp);
 
   return 0;
 }