Regenerate MLIR Bindings (#1936)

Co-authored-by: enzyme-ci-bot[bot] <78882869+enzyme-ci-bot[bot]@users.noreply.github.com>

Author: github-actions[bot]

src/mlir/Dialects/Arith.jl

@@ -548,12 +548,18 @@ is undefined behavior. When applied to `vector` and `tensor` values, the
 behavior is undefined if _any_ of its elements are divided by zero or has a
 signed division overflow.
 
+If the `exact` attribute is present, the result value is poison if `lhs` is
+not a multiple of `rhs`.
+
 # Example
 
 ```mlir
 // Scalar signed integer division.
 %a = arith.divsi %b, %c : i64
 
+// Scalar signed integer division where %b is known to be a multiple of %c.
+%a = arith.divsi %b, %c exact : i64
+
 // SIMD vector element-wise division.
 %f = arith.divsi %g, %h : vector<4xi32>
 
@@ -562,14 +568,19 @@ signed division overflow.
 ```
 """
 function divsi(
-    lhs::Value, rhs::Value; result=nothing::Union{Nothing,IR.Type}, location=Location()
+    lhs::Value,
+    rhs::Value;
+    result=nothing::Union{Nothing,IR.Type},
+    isExact=nothing,
+    location=Location(),
 )
     op_ty_results = IR.Type[]
     operands = Value[lhs, rhs]
     owned_regions = Region[]
     successors = Block[]
     attributes = NamedAttribute[]
     !isnothing(result) && push!(op_ty_results, result)
+    !isnothing(isExact) && push!(attributes, namedattribute("isExact", isExact))
 
     return create_operation(
         "arith.divsi",
@@ -594,12 +605,18 @@ Division by zero is undefined behavior. When applied to `vector` and
 `tensor` values, the behavior is undefined if _any_ elements are divided by
 zero.
 
+If the `exact` attribute is present, the result value is poison if `lhs` is
+not a multiple of `rhs`.
+
 # Example
 
 ```mlir
 // Scalar unsigned integer division.
 %a = arith.divui %b, %c : i64
 
+// Scalar unsigned integer division where %b is known to be a multiple of %c.
+%a = arith.divui %b, %c exact : i64
+
 // SIMD vector element-wise division.
 %f = arith.divui %g, %h : vector<4xi32>
 
@@ -608,14 +625,19 @@ zero.
 ```
 """
 function divui(
-    lhs::Value, rhs::Value; result=nothing::Union{Nothing,IR.Type}, location=Location()
+    lhs::Value,
+    rhs::Value;
+    result=nothing::Union{Nothing,IR.Type},
+    isExact=nothing,
+    location=Location(),
 )
     op_ty_results = IR.Type[]
     operands = Value[lhs, rhs]
     owned_regions = Region[]
     successors = Block[]
     attributes = NamedAttribute[]
     !isnothing(result) && push!(op_ty_results, result)
+    !isnothing(isExact) && push!(attributes, namedattribute("isExact", isExact))
 
     return create_operation(
         "arith.divui",
@@ -1789,25 +1811,33 @@ output are filled with copies of the most-significant bit of the shifted value
 operand is greater or equal than bitwidth of the first operand, then the operation
 returns poison.
 
+If the `exact` attribute is present, the result value of shrsi is a poison
+value if any of the bits shifted out are non-zero.
+
 # Example
 
 ```mlir
-%1 = arith.constant 160 : i8               // %1 is 0b10100000
+%1 = arith.constant 160 : i8         // %1 is 0b10100000
 %2 = arith.constant 3 : i8
-%3 = arith.shrsi %1, %2 : (i8, i8) -> i8   // %3 is 0b11110100
-%4 = arith.constant 96 : i8                   // %4 is 0b01100000
-%5 = arith.shrsi %4, %2 : (i8, i8) -> i8   // %5 is 0b00001100
+%3 = arith.shrsi %1, %2 exact : i8   // %3 is 0b11110100
+%4 = arith.constant 98 : i8          // %4 is 0b01100010
+%5 = arith.shrsi %4, %2 : i8         // %5 is 0b00001100
 ```
 """
 function shrsi(
-    lhs::Value, rhs::Value; result=nothing::Union{Nothing,IR.Type}, location=Location()
+    lhs::Value,
+    rhs::Value;
+    result=nothing::Union{Nothing,IR.Type},
+    isExact=nothing,
+    location=Location(),
 )
     op_ty_results = IR.Type[]
     operands = Value[lhs, rhs]
     owned_regions = Region[]
     successors = Block[]
     attributes = NamedAttribute[]
     !isnothing(result) && push!(op_ty_results, result)
+    !isnothing(isExact) && push!(attributes, namedattribute("isExact", isExact))
 
     return create_operation(
         "arith.shrsi",
@@ -1830,23 +1860,33 @@ and the second operand is interpreted as unsigned. The high order bits are alway
 filled with zeros. If the value of the second operand is greater or equal than the
 bitwidth of the first operand, then the operation returns poison.
 
+If the `exact` attribute is present, the result value of shrui is a poison
+value if any of the bits shifted out are non-zero.
+
 # Example
 
 ```mlir
-%1 = arith.constant 160 : i8               // %1 is 0b10100000
+%1 = arith.constant 160 : i8        // %1 is 0b10100000
 %2 = arith.constant 3 : i8
-%3 = arith.shrui %1, %2 : (i8, i8) -> i8   // %3 is 0b00010100
+%3 = arith.constant 6 : i8
+%4 = arith.shrui %1, %2 exact : i8  // %4 is 0b00010100
+%5 = arith.shrui %1, %3 : i8        // %3 is 0b00000010
 ```
 """
 function shrui(
-    lhs::Value, rhs::Value; result=nothing::Union{Nothing,IR.Type}, location=Location()
+    lhs::Value,
+    rhs::Value;
+    result=nothing::Union{Nothing,IR.Type},
+    isExact=nothing,
+    location=Location(),
 )
     op_ty_results = IR.Type[]
     operands = Value[lhs, rhs]
     owned_regions = Region[]
     successors = Block[]
     attributes = NamedAttribute[]
     !isnothing(result) && push!(op_ty_results, result)
+    !isnothing(isExact) && push!(attributes, namedattribute("isExact", isExact))
 
     return create_operation(
         "arith.shrui",

src/mlir/Dialects/MosaicGPU.jl

@@ -215,6 +215,10 @@ does not.
 
 The `predicate` allows scheduling the transfer conditionally. The async copy
 is always scheduled by at most a single lane in the warpgroup.
+
+The `reduction_op` attribute can be provided to perform a reduction when
+storing to GMEM. For example, using `add` will add the SMEM values to
+existing values in GMEM.
 """
 function async_store(
     source::Value,
@@ -223,6 +227,7 @@ function async_store(
     predicate=nothing::Union{Nothing,Value};
     slice_lengths,
     commit_group=nothing,
+    reduction_op=nothing,
     location=Location(),
 )
     op_ty_results = IR.Type[]
@@ -237,6 +242,8 @@ function async_store(
     )
     !isnothing(commit_group) &&
         push!(attributes, namedattribute("commit_group", commit_group))
+    !isnothing(reduction_op) &&
+        push!(attributes, namedattribute("reduction_op", reduction_op))
 
     return create_operation(
         "mosaic_gpu.async_store",
@@ -639,7 +646,7 @@ The `smem_ptr` is a pointer in SMEM where a pointer to the allocated
 TMEM will be stored. The op returns a memref to the allocated TMEM. The
 result must have a shape with dimensions [rows, logical_columns]. If
 `packing` is 1, then the number of logical (unpacked) columns is equal to
-the number of allocated columns in TMEM. Otherwise, these equations
+the number of allocated columns in TMEM. Otherwise, these constraints
 must hold:
 
     packing = 32 / bitwidth(element type of result)