The affine.apply operation applies an affine mapping to a list of SSA values, yielding a single SSA value. The number of dimension and symbol arguments to affine.apply must be equal to the respective number of dimensional and symbolic inputs to the affine mapping; the affine mapping has to be one-dimensional, and so the affine.apply operation always returns one value. The input operands and result must all have ‘index’ type.
The affine.for operation represents an affine loop nest. It has one region containing its body. This region must contain one block that terminates with affine.yield. Note: when affine.for is printed in custom format, the terminator is omitted. The block has one argument of index type that represents the induction variable of the loop.
The affine.for operation executes its body a number of times iterating from a lower bound to an upper bound by a stride. The stride, represented by step, is a positive constant integer which defaults to "1" if not present. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound.
The lower and upper bounds of a affine.for operation are represented as an application of an affine mapping to a list of SSA values passed to the map. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols.
The affine mappings for the bounds may return multiple results, in which case the max/min keywords are required (for the lower/upper bound respectively), and the bound is the maximum/minimum of the returned values. There is no semantic ambiguity, but MLIR syntax requires the use of these keywords to make things more obvious to human readers.
Many upper and lower bounds are simple, so MLIR accepts two custom form syntaxes: the form that accepts a single 'ssa-id' (e.g. %N) is shorthand for applying that SSA value to a function that maps a single symbol to itself, e.g., ()[s]->(s)()[%N]. The integer literal form (e.g. -42) is shorthand for a nullary mapping function that returns the constant value (e.g. ()->(-42)()).
Example showing reverse iteration of the inner loop:
affine.for can also operate on loop-carried variables (iter_args) and return the final values after loop termination. The initial values of the variables are passed as additional SSA operands to the affine.for following the operands for the loop's lower and upper bounds. The operation's region has equivalent arguments for each variable representing the value of the variable at the current iteration.
The region must terminate with an affine.yield that passes all the current iteration variables to the next iteration, or to the affine.for's results if at the last iteration. For affine.for's that execute zero iterations, the 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.func @reduce(%buffer: memref<1024xf32>) -> (f32) {
+ // Initial sum set to 0.
+ %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
+ }
+ return %sum : f32
+}
If the affine.for defines any values, a yield terminator must be explicitly present. The number and types of the "affine.for" results must match the initial values in the iter_args binding and the yield operands.
The affine.if operation restricts execution to a subset of the loop iteration space defined by an integer set (a conjunction of affine constraints). A single affine.if may end with an optional else clause.
The condition of the affine.if is represented by an integer set (a conjunction of affine constraints), and the SSA values bound to the dimensions and symbols in the integer set. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols.
The affine.if operation contains two regions for the "then" and "else" clauses. affine.if may return results that are defined in its regions. The values defined are determined by which execution path is taken. Each region of the affine.if must contain a single block with no arguments, and be terminated by affine.yield. If affine.if defines no values, the affine.yield can be left out, and will be inserted implicitly. Otherwise, it must be explicit. If no values are defined, the else block may be empty (i.e. contain no blocks).
The affine.load op reads an element from a memref, where the index for each memref dimension is an affine expression of loop induction variables and symbols. The output of affine.load is a new value with the same type as the elements of the memref. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.min operation applies an affine mapping to a list of SSA values, and returns the minimum value of all result expressions. The number of dimension and symbol arguments to affine.min must be equal to the respective number of dimensional and symbolic inputs to the affine mapping; the affine.min operation always returns one value. The input operands and result must all have 'index' type.
The affine.parallel operation represents a hyper-rectangular affine parallel band, defining zero or more SSA values for its induction variables. It has one region capturing the parallel band body. The induction variables are represented as arguments of this region. These SSA values always have type index, which is the size of the machine word. The strides, represented by steps, are positive constant integers which defaults to "1" if not present. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound. The body region must contain exactly one block that terminates with affine.yield.
The lower and upper bounds of a parallel operation are represented as an application of an affine mapping to a list of SSA values passed to the map. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols. The list of expressions in each map is interpreted according to the respective bounds group attribute. If a single expression belongs to the group, then the result of this expression is taken as a lower(upper) bound of the corresponding loop induction variable. If multiple expressions belong to the group, then the lower(upper) bound is the max(min) of these values obtained from these expressions. The loop band has as many loops as elements in the group bounds attributes.
Each value yielded by affine.yield will be accumulated/reduced via one of the reduction methods defined in the AtomicRMWKind enum. The order of reduction is unspecified, and lowering may produce any valid ordering. Loops with a 0 trip count will produce as a result the identity value associated with each reduction (i.e. 0.0 for addf, 1.0 for mulf). Assign reductions for loops with a trip count != 1 produces undefined results.
Note: Calling AffineParallelOp::build will create the required region and block, and insert the required terminator if it is trivial (i.e. no values are yielded). Parsing will also create the required region, block, and terminator, even when they are missing from the textual representation.
The affine.prefetch op prefetches data from a memref location described with an affine subscript similar to affine.load, and has three attributes: a read/write specifier, a locality hint, and a cache type specifier as shown below:
mlir
affine.prefetch %0[%i, %j + 5], read, locality<3>, data : memref<400x400xi32>
The read/write specifier is either 'read' or 'write', the locality hint specifier ranges from locality<0> (no locality) to locality<3> (extremely local keep in cache). The cache type specifier is either 'data' or 'instr' and specifies whether the prefetch is performed on data cache or on instruction cache.
The affine.store op writes an element to a memref, where the index for each memref dimension is an affine expression of loop induction variables and symbols. The affine.store op stores a new value which is the same type as the elements of the memref. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.vector_load is the vector counterpart of affine.load. It reads a slice from a MemRef, supplied as its first operand, into a vector of the same base elemental type. The index for each memref dimension is an affine expression of loop induction variables and symbols. These indices determine the start position of the read within the memref. The shape of the return vector type determines the shape of the slice read from the memref. This slice is contiguous along the respective dimensions of the shape. Strided vector loads will be supported in the future. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.vector_store is the vector counterpart of affine.store. It writes a vector, supplied as its first operand, into a slice within a MemRef of the same base elemental type, supplied as its second operand. The index for each memref dimension is an affine expression of loop induction variables and symbols. These indices determine the start position of the write within the memref. The shape of th input vector determines the shape of the slice written to the memref. This slice is contiguous along the respective dimensions of the shape. Strided vector stores will be supported in the future. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.yield yields zero or more SSA values from an affine op region and terminates the region. The semantics of how the values yielded are used is defined by the parent operation. If affine.yield has any operands, the operands must match the parent operation's results. If the parent operation defines no values, then the affine.yield may be left out in the custom syntax and the builders will insert one implicitly. Otherwise, it has to be present in the syntax to indicate which values are yielded.
Converts certain expressions like control flow into a Reactant friendly form. Importantly, if no traced value is found inside the expression, then there is no overhead.
Currently Supported
if conditions (with elseif and other niceties) (@trace if ...)
if statements with a preceeding assignment (@trace a = if ...) (note the positioning of the macro needs to be before the assignment and not before the if)
for statements with a single induction variable iterating over a syntactic StepRange of integers.
Special Considerations
Apply @trace only at the outermost if. Nested if statements will be automatically expanded into the correct form.
Extended Help
Caveats (Deviations from Core Julia Semantics)
New variables introduced
julia
@trace if x > 0
+ y = x + 1
+ p = 1
+else
+ y = x - 1
+end
In the outer scope p is not defined if x ≤ 0. However, for the traced version, it is defined and set to a dummy value.
Short Circuiting Operations
julia
@trace if x > 0 && z > 0
+ y = x + 1
+else
+ y = x - 1
+end
&& and || are short circuiting operations. In the traced version, we replace them with & and | respectively.
Type-Unstable Branches
julia
@trace if x > 0
+ y = 1.0f0
+else
+ y = 1.0
+end
This will not compile since y is a Float32 in one branch and a Float64 in the other. You need to ensure that all branches have the same type.
Another example is the following for loop which changes the type of x between iterations.
julia
x = ... # ConcreteRArray{Int64, 1}
+for i in 1f0:0.5f0:10f0
+ x = x .+ i # ConcreteRArray{Float32, 1}
+end
Certain Symbols are Reserved
Symbols like [😦😃, :nothing, :missing, :Inf, :Inf16, :Inf32, :Inf64, :Base, :Core] are not allowed as variables in @trace expressions. While certain cases might work but these are not guaranteed to work. For example, the following will not work:
julia
function fn(x)
+ nothing = sum(x)
+ @trace if nothing > 0
+ y = 1.0
+ else
+ y = 2.0
+ end
+ return y, nothing
+end
Generate Julia code to wrap the XLA buffers back into the output result datatypes. The name is due to its similarity to the unflatten function in jax.tree_util.register_pytree_node.
Generate Julia code to extract the XLA buffers from input arguments. The name is due to its similarity to the flatten function in jax.tree_util.register_pytree_node.
Arguments
linear_args: A list of arguments to be flattened.
Returns
flatten_names: A list of Symbols representing the names of the flattened arguments.
flatten_code: A list of Exprs to extract the XLA buffers from the input arguments.
Note
The linearized arguments do not directly refer to the are the arguments that have been flattened into a single list.
The addf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit addition on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the addition modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The addi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Performs (N+1)-bit addition on zero-extended operands. Returns two results: the N-bit sum (same type as both operands), and the overflow bit (boolean-like), where 1 indicates unsigned addition overflow, while 0 indicates no overflow.
The andi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
Bitcast an integer or floating point value to an integer or floating point value of equal bit width. When operating on vectors, casts elementwise.
Note that this implements a logical bitcast independent of target endianness. This allows constant folding without target information and is consitent with the bitcast constant folders in LLVM (see https://github.com/llvm/llvm-project/blob/18c19414eb/llvm/lib/IR/ConstantFold.cpp#L168) For targets where the source and target type have the same endianness (which is the standard), this cast will also change no bits at runtime, but it may still require an operation, for example if the machine has different floating point and integer register files. For targets that have a different endianness for the source and target types (e.g. float is big-endian and integer is little-endian) a proper lowering would add operations to swap the order of words in addition to the bitcast.
Signed integer division. Rounds towards positive infinity, i.e. 7 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Unsigned integer division. Rounds towards positive infinity. Treats the leading bit as the most significant, i.e. for i16 given two's complement representation, 6 / -2 = 6 / (2^16 - 2) = 1.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
The cmpf operation compares its two operands according to the float comparison rules and the predicate specified by the respective attribute. The predicate defines the type of comparison: (un)orderedness, (in)equality and signed less/greater than (or equal to) as well as predicates that are always true or false. The operands must have the same type, and this type must be a float type, or a vector or tensor thereof. The result is an i1, or a vector/tensor thereof having the same shape as the inputs. Unlike cmpi, the operands are always treated as signed. The u prefix indicates unordered comparison, not unsigned comparison, so "une" means unordered or not equal. For the sake of readability by humans, custom assembly form for the operation uses a string-typed attribute for the predicate. The value of this attribute corresponds to lower-cased name of the predicate constant, e.g., "one" means "ordered not equal". The string representation of the attribute is merely a syntactic sugar and is converted to an integer attribute by the parser.
The cmpi operation is a generic comparison for integer-like types. Its two arguments can be integers, vectors or tensors thereof as long as their types match. The operation produces an i1 for the former case, a vector or a tensor of i1 with the same shape as inputs in the other cases.
Its first argument is an attribute that defines which type of comparison is performed. The following comparisons are supported:
equal (mnemonic: "eq"; integer value: 0)
not equal (mnemonic: "ne"; integer value: 1)
signed less than (mnemonic: "slt"; integer value: 2)
signed less than or equal (mnemonic: "sle"; integer value: 3)
signed greater than (mnemonic: "sgt"; integer value: 4)
signed greater than or equal (mnemonic: "sge"; integer value: 5)
unsigned less than (mnemonic: "ult"; integer value: 6)
unsigned less than or equal (mnemonic: "ule"; integer value: 7)
unsigned greater than (mnemonic: "ugt"; integer value: 8)
unsigned greater than or equal (mnemonic: "uge"; integer value: 9)
The result is 1 if the comparison is true and 0 otherwise. For vector or tensor operands, the comparison is performed elementwise and the element of the result indicates whether the comparison is true for the operand elements with the same indices as those of the result.
Note: while the custom assembly form uses strings, the actual underlying attribute has integer type (or rather enum class in C++ code) as seen from the generic assembly form. String literals are used to improve readability of the IR by humans.
This operation only applies to integer-like operands, but not floats. The main reason being that comparison operations have diverging sets of attributes: integers require sign specification while floats require various floating point-related particularities, e.g., -ffast-math behavior, IEEE754 compliance, etc (rationale). The type of comparison is specified as attribute to avoid introducing ten similar operations, taking into account that they are often implemented using the same operation downstream (rationale). The separation between signed and unsigned order comparisons is necessary because of integers being signless. The comparison operation must know how to interpret values with the foremost bit being set: negatives in two's complement or large positives (rationale).
Example
mlir
// Custom form of scalar "signed less than" comparison.
+%x = arith.cmpi slt, %lhs, %rhs : i32
+
+// Generic form of the same operation.
+%x = "arith.cmpi"(%lhs, %rhs) {predicate = 2 : i64} : (i32, i32) -> i1
+
+// Custom form of vector equality comparison.
+%x = arith.cmpi eq, %lhs, %rhs : vector<4xi64>
+
+// Generic form of the same operation.
+%x = "arith.cmpi"(%lhs, %rhs) {predicate = 0 : i64}
+ : (vector<4xi64>, vector<4xi64>) -> vector<4xi1>
The constant operation produces an SSA value equal to some integer or floating-point constant specified by an attribute. This is the way MLIR forms simple integer and floating point constants.
Signed integer division. Rounds towards zero. Treats the leading bit as sign, i.e. 6 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Unsigned integer division. Rounds towards zero. Treats the leading bit as the most significant, i.e. for i16 given two's complement representation, 6 / -2 = 6 / (2^16 - 2) = 0.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
Cast a floating-point value to a larger floating-point-typed value. The destination type must to be strictly wider than the source type. When operating on vectors, casts elementwise.
The integer sign extension operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be larger than the input bit-width (N > M). The top-most (N - M) bits of the output are filled with copies of the most-significant bit of the input.
Example
mlir
%1 = arith.constant 5 : i3 // %1 is 0b101
+%2 = arith.extsi %1 : i3 to i6 // %2 is 0b111101
+%3 = arith.constant 2 : i3 // %3 is 0b010
+%4 = arith.extsi %3 : i3 to i6 // %4 is 0b000010
+
+%5 = arith.extsi %0 : vector<2 x i32> to vector<2 x i64>
The integer zero extension operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be larger than the input bit-width (N > M). The top-most (N - M) bits of the output are filled with zeros.
Example
mlir
%1 = arith.constant 5 : i3 // %1 is 0b101
+ %2 = arith.extui %1 : i3 to i6 // %2 is 0b000101
+ %3 = arith.constant 2 : i3 // %3 is 0b010
+ %4 = arith.extui %3 : i3 to i6 // %4 is 0b000010
+
+ %5 = arith.extui %0 : vector<2 x i32> to vector<2 x i64>
Signed integer division. Rounds towards negative infinity, i.e. 5 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Cast from a value interpreted as floating-point to the nearest (rounding towards zero) signed integer value. When operating on vectors, casts elementwise.
Cast from a value interpreted as floating-point to the nearest (rounding towards zero) unsigned integer value. When operating on vectors, casts elementwise.
Casts between scalar or vector integers and corresponding 'index' scalar or vectors. Index is an integer of platform-specific bit width. If casting to a wider integer, the value is sign-extended. If casting to a narrower integer, the value is truncated.
Casts between scalar or vector integers and corresponding 'index' scalar or vectors. Index is an integer of platform-specific bit width. If casting to a wider integer, the value is zero-extended. If casting to a narrower integer, the value is truncated.
Returns the maximum of the two arguments. If the arguments are -0.0 and +0.0, then the result is either of them. If one of the arguments is NaN, then the result is the other argument.
Returns the minimum of the two arguments. If the arguments are -0.0 and +0.0, then the result is either of them. If one of the arguments is NaN, then the result is the other argument.
The mulf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit multiplication on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the multiplication modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The muli operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Performs (2*N)-bit multiplication on sign-extended operands. Returns two N-bit results: the low and the high halves of the product. The low half has the same value as the result of regular multiplication arith.muli with the same operands.
Performs (2*N)-bit multiplication on zero-extended operands. Returns two N-bit results: the low and the high halves of the product. The low half has the same value as the result of regular multiplication arith.muli with the same operands.
The negf operation computes the negation of a given value. It takes one operand and returns one result of the same type. This type may be a float scalar type, a vector whose element type is float, or a tensor of floats. It has no standard attributes.
The ori operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
The arith.select operation chooses one value based on a binary condition supplied as its first operand.
If the value of the first operand (the condition) is 1, then the second operand is returned, and the third operand is ignored, even if it was poison.
If the value of the first operand (the condition) is 0, then the third operand is returned, and the second operand is ignored, even if it was poison.
If the value of the first operand (the condition) is poison, then the operation returns poison.
The operation applies to vectors and tensors elementwise given the shape of all operands is identical. The choice is made for each element individually based on the value at the same position as the element in the condition operand. If an i1 is provided as the condition, the entire vector or tensor is chosen.
Example
mlir
// Custom form of scalar selection.
+%x = arith.select %cond, %true, %false : i32
+
+// Generic form of the same operation.
+%x = "arith.select"(%cond, %true, %false) : (i1, i32, i32) -> i32
+
+// Element-wise vector selection.
+%vx = arith.select %vcond, %vtrue, %vfalse : vector<42xi1>, vector<42xf32>
+
+// Full vector selection.
+%vx = arith.select %cond, %vtrue, %vfalse : vector<42xf32>
The shli operation shifts the integer value of the first operand to the left by the integer value of the second operand. The second operand is interpreted as unsigned. The low order bits are 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.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
The shrsi operation shifts an integer value of the first operand to the right by the value of the second operand. The first operand is interpreted as signed, and the second operand is interpreter as unsigned. The high order bits in the output are filled with copies of the most-significant bit of the shifted value (which means that the sign of the value is preserved). If the value of the second operand is greater or equal than bitwidth of the first operand, then the operation returns poison.
The shrui operation shifts an integer value of the first operand to the right by the value of the second operand. The first operand is interpreted as unsigned, and the second operand is interpreted as unsigned. The high order bits are always 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.
Cast from a value interpreted as a signed integer to the corresponding floating-point value. If the value cannot be exactly represented, it is rounded using the default rounding mode. When operating on vectors, casts elementwise.
The subf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit subtraction on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the subtraction modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The subi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Truncate a floating-point value to a smaller floating-point-typed value. The destination type must be strictly narrower than the source type. If the value cannot be exactly represented, it is rounded using the provided rounding mode or the default one if no rounding mode is provided. When operating on vectors, casts elementwise.
The integer truncation operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be smaller than the input bit-width (N < M). The top-most (N - M) bits of the input are discarded.
Example
mlir
%1 = arith.constant 21 : i5 // %1 is 0b10101
+ %2 = arith.trunci %1 : i5 to i4 // %2 is 0b0101
+ %3 = arith.trunci %1 : i5 to i3 // %3 is 0b101
+
+ %5 = arith.trunci %0 : vector<2 x i32> to vector<2 x i16>
Cast from a value interpreted as unsigned integer to the corresponding floating-point value. If the value cannot be exactly represented, it is rounded using the default rounding mode. When operating on vectors, casts elementwise.
The xori operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
A module represents a top-level container operation. It contains a single graph region containing a single block which can contain any operations and does not have a terminator. Operations within this region cannot implicitly capture values defined outside the module, i.e. Modules are IsolatedFromAbove. Modules have an optional symbol name which can be used to refer to them in operations.
An unrealized_conversion_cast operation represents an unrealized conversion from one set of types to another, that is used to enable the inter-mixing of different type systems. This operation should not be attributed any special representational or execution semantics, and is generally only intended to be used to satisfy the temporary intermixing of type systems during the conversion of one type system to another.
This operation may produce results of arity 1-N, and accept as input operands of arity 0-N.
Example
mlir
// An unrealized 0-1 conversion. These types of conversions are useful in
+// cases where a type is removed from the type system, but not all uses have
+// been converted. For example, imagine we have a tuple type that is
+// expanded to its element types. If only some uses of an empty tuple type
+// instance are converted we still need an instance of the tuple type, but
+// have no inputs to the unrealized conversion.
+%result = unrealized_conversion_cast to !bar.tuple_type<>
+
+// An unrealized 1-1 conversion.
+%result1 = unrealized_conversion_cast %operand : !foo.type to !bar.lowered_type
+
+// An unrealized 1-N conversion.
+%results2:2 = unrealized_conversion_cast %tuple_operand : !foo.tuple_type<!foo.type, !foo.type> to !foo.type, !foo.type
+
+// An unrealized N-1 conversion.
+%result3 = unrealized_conversion_cast %operand, %operand : !foo.type, !foo.type to !bar.tuple_type<!foo.type, !foo.type>
Compares lhs and rhs elementwise according to comparison_direction and compare_type. If unspecified, compare_type is FLOAT for float element types, SIGNED for signed element types and UNSIGNED for unsigned element types.
If the input is a vector (rank-1), finds the k largest entries in the vector and outputs their values and indices as vectors. Thus values[j] is the j-th largest entry in input, and its index is indices[j].
For matrices (resp. higher rank input), computes the top k entries in each row (resp. vector along the last dimension). Thus,
The func.call operation represents a direct call to a function that is within the same symbol scope as the call. The operands and result types of the call must match the specified function type. The callee is encoded as a symbol reference attribute named "callee".
The func.call_indirect operation represents an indirect call to a value of function type. The operands and result types of the call must match the specified function type.
The func.constant operation produces an SSA value from a symbol reference to a func.func operation
Example
mlir
// Reference to function @myfn.
+%2 = func.constant @myfn : (tensor<16xf32>, f32) -> tensor<16xf32>
+
+// Equivalent generic forms
+%2 = "func.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).
Operations within the function cannot implicitly capture values defined outside of the function, i.e. Functions are IsolatedFromAbove. All external references must use function arguments or attributes that establish a symbolic connection (e.g. symbols referenced by name via a string attribute like SymbolRefAttr). An external function declaration (used when referring to a function declared in some other module) has no body. While the MLIR textual form provides a nice inline syntax for function arguments, they are internally represented as “block arguments” to the first block in the region.
Only dialect attribute names may be specified in the attribute dictionaries for function arguments, results, or the function itself.
Example
mlir
// External function definitions.
+func.func private @abort()
+func.func private @scribble(i32, i64, memref<? x 128 x f32, #layout_map0>) -> f64
+
+// A function that returns its argument twice:
+func.func @count(%x: i64) -> (i64, i64)
+ attributes {fruit = "banana"} {
+ return %x, %x: i64, i64
+}
+
+// A function with an argument attribute
+func.func private @example_fn_arg(%x: i32 {swift.self = unit})
+
+// A function with a result attribute
+func.func private @example_fn_result() -> (f64 {dialectName.attrName = 0 : i64})
+
+// A function with an attribute
+func.func private @example_fn_attr() attributes {dialectName.attrName = false}
The func.return operation represents a return operation within a function. The operation takes variable number of operands and produces no results. The operand number and types must match the signature of the function that contains the operation.
Creates an affine map with results defined by the given list of affine expressions. The map resulting map also has the requested number of input dimensions and symbols, regardless of them being used in the results.
Creates a floating point attribute in the given context with the given double value and double-precision FP semantics. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an ExecutionEngine for the provided ModuleOp. The ModuleOp is expected to be "translatable" to LLVM IR (only contains operations in dialects that implement the LLVMTranslationDialectInterface). The module ownership stays with the client and can be destroyed as soon as the call returns. optLevel is the optimization level to be used for transformation and code generation. LLVM passes at optLevel are run before code generation. The number and array of paths corresponding to shared libraries that will be loaded are specified via numPaths and sharedLibPaths respectively. TODO: figure out other options.
Gets or creates a new integer set in the given context. The set is defined by a list of affine constraints, with the given number of input dimensions and symbols, which are treated as either equalities (eqflags is 1) or inequalities (eqflags is 0). Both constraints and eqflags need to be arrays of the same length.
A logical result value, essentially a boolean with named states. LLVM convention for using boolean values to designate success or failure of an operation is a moving target, so MLIR opted for an explicit class. Instances of LogicalResult must only be inspected using the associated functions.
Nest an OpPassManager under the provided OpPassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed.
Nest an OpPassManager under the top-level PassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed. To further nest more OpPassManager under the newly returned one, see mlirOpPassManagerNest below.
Checks if two integer set objects are equal. This is a "shallow" comparison of two objects. Only the sets with some small number of constraints are uniqued and compare equal here. Set objects that represent the same integer set with different constraints may be considered non-equal by this check. Set difference followed by an (expensive) emptiness check should be used to check equivalence of the underlying integer sets.
Takes an operation owned by the caller and inserts it as index to the block. This is an expensive operation that scans the block linearly, prefer insertBefore/After instead.
Takes a block owned by the caller and inserts it at index to the given region. This is an expensive operation that linearly scans the region, prefer insertAfter/Before instead.
Inserts the given operation into the given symbol table. The operation must have the symbol trait. If the symbol table already has a symbol with the same name, renames the symbol being inserted to ensure name uniqueness. Note that this does not move the operation itself into the block of the symbol table operation, this should be done separately. Returns the name of the symbol after insertion.
Apply AffineExpr::replace(map) to each of the results and return a new new AffineMap with the new results and the specified number of dims and symbols.
Gets or creates a new integer set in which the values and dimensions of the given set are replaced with the given affine expressions. dimReplacements and symbolReplacements are expected to point to at least as many consecutive expressions as the given set has dimensions and symbols, respectively. The new set will have numResultDims and numResultSymbols dimensions and symbols, respectively.
Creates a dense elements attribute that has the same data as the given dense elements attribute and a different shaped type. The new type must have the same total number of elements.
Creates a dense elements attribute with the given shaped type from elements of a specific type. Expects the element type of the shaped type to match the data element type.
Creates a MemRef type with the given rank and shape, a potentially empty list of affine layout maps, the given memory space and element type, in the same context as element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates a MemRef type with the given rank, shape, memory space and element type in the same context as the element type. The type has no affine maps, i.e. represents a default row-major contiguous memref. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an Unranked MemRef type with the given element type and in the given memory space. The type is owned by the context of element type. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an identity affine map on the most minor dimensions in the context. The affine map is owned by the context. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an opaque attribute in the given context associated with the dialect identified by its namespace. The attribute contains opaque byte data of the specified length (data need not be null-terminated).
Creates an opaque type in the given context associated with the dialect identified by its namespace. The type contains opaque byte data of the specified length (data need not be null-terminated).
Creates an affine map with a permutation expression and its size in the context. The permutation expression is a non-empty vector of integers. The elements of the permutation vector must be continuous from 0 and cannot be repeated (i.e. [1,2,0] is a valid permutation. [2,0] or [1,1,2] is an invalid invalid permutation). The affine map is owned by the context.
Creates a symbol reference attribute in the given context referencing a symbol identified by the given string inside a list of nested references. Each of the references in the list must not be nested.
Creates a tensor type of a fixed rank with the given shape, element type, and optional encoding in the same context as the element type. The type is owned by the context. Tensor types without any specific encoding field should assign mlirAttributeGetNull to this parameter. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an unranked tensor type with the given element type in the same context as the element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates a vector type of the shape identified by its rank and dimensions, with the given element type in the same context as the element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Add a pass and transfer ownership to the provided OpPassManager. If the pass is not a generic operation pass or matching the type of the provided OpPassManager, a new OpPassManager is implicitly nested under the provided OpPassManager.
Add a pass and transfer ownership to the provided top-level PassManager. If the pass is not a generic operation pass or a ModulePass, a new OpPassManager is implicitly nested under the provided PassManager.
Parse a sequence of textual MLIR pass pipeline elements and add them to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Takes an operation owned by the caller and inserts it after the (non-owned) reference operation in the given block. If the reference is null, prepends the operation. Otherwise, the reference must belong to the block.
Takes a block owned by the caller and inserts it after the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, prepends the block to the region.
Takes an operation owned by the caller and inserts it before the (non-owned) reference operation in the given block. If the reference is null, appends the operation. Otherwise, the reference must belong to the block.
Takes a block owned by the caller and inserts it before the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, appends the block to the region.
Returns whether the given fully-qualified operation (i.e. 'dialect.operation') is registered with the context. This will return true if the dialect is loaded and the operation is registered within the dialect.
Checks whether the given affine map is an identity affine map. The function asserts that the number of dimensions is greater or equal to the number of results.
Looks up a symbol with the given name in the given symbol table and returns the operation that corresponds to the symbol. If the symbol cannot be found, returns a null operation.
Returns the affine map consisting of the most major nresults results. Returns the null AffineMap if the nresults is equal to zero. Returns the affineMap if nresults is greater or equals to number of results of the given affine map.
Returns the affine map consisting of the most minor nresults results. Returns the null AffineMap if the nresults is equal to zero. Returns the affineMap if nresults is greater or equals to number of results of the given affine map.
Moves the given operation immediately after the other operation in its parent block. The given operation may be owned by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Moves the given operation immediately before the other operation in its parent block. The given operation may be owner by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Interface used to provide a module to JIT or interpreter. This is now just a synonym for llvm::Module, but we have to keep using the different type to keep binary compatibility.
Diagnostic handler type. Accepts a reference to a diagnostic, which is only guaranteed to be live during the call. The handler is passed the userData that was provided when the handler was attached to a context. If the handler processed the diagnostic completely, it is expected to return success. Otherwise, it is expected to return failure to indicate that other handlers should attempt to process the diagnostic.
Structure of external MlirPass callbacks. All callbacks are required to be set unless otherwise specified.
Field
Note
construct
This callback is called from the pass is created. This is analogous to a C++ pass constructor.
destruct
This callback is called when the pass is destroyed This is analogous to a C++ pass destructor.
initialize
This callback is optional. The callback is called before the pass is run, allowing a chance to initialize any complex state necessary for running the pass. See Pass::initialize(MLIRContext *).
clone
This callback is called when the pass is cloned. See Pass::clonePass().
run
This callback is called when the pass is run. See Pass::runOnOperation().
A logical result value, essentially a boolean with named states. LLVM convention for using boolean values to designate success or failure of an operation is a moving target, so MLIR opted for an explicit class. Instances of MlirLogicalResult must only be inspected using the associated functions.
This class contains all the information necessary to construct the operation. It owns the MlirRegions it has pointers to and does not own anything else. By default, the state can be constructed from a name and location, the latter being also used to access the context, and has no other components. These components can be added progressively until the operation is constructed. Users are not expected to rely on the internals of this class and should use mlirOperationState* functions instead.
These callbacks are used to return multiple shaped type components from functions while transferring ownership to the caller. The first argument is the has rank boolean followed by the the rank and a pointer to the shape (if applicable). The next argument is the element type, then the attribute. The last argument is an opaque pointer forwarded to the callback by the caller. This callback will be called potentially multiple times for each shaped type components.
This function is called back by the functions that need to return a reference to the portion of the string with the following arguments: - an MlirStringRef representing the current portion of the string - a pointer to user data forwarded from the printing call.
These callbacks are used to return multiple types from functions while transferring ownership to the caller. The first argument is the number of consecutive elements pointed to by the second argument. The third argument is an opaque pointer forwarded to the callback by the caller.
This function parses the given arguments using the LLVM command line parser. Note that the only stable thing about this function is its signature; you cannot rely on any particular set of command line arguments being interpreted the same way across LLVM versions.
This function will search through all previously loaded dynamic libraries for the symbol symbolName. If it is found, the address of that symbol is returned. If not, null is returned.
Prints an affine expression by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Returns the simplified affine map resulting from dropping the symbols that do not appear in any of the individual maps in affineMaps. Asserts that all maps in affineMaps are normalized to the same number of dims and symbols. Takes a callback populateResult to fill the res container with value m at entry idx. This allows returning without worrying about ownership considerations.
Creates an affine map with results defined by the given list of affine expressions. The map resulting map also has the requested number of input dimensions and symbols, regardless of them being used in the results.
Returns the affine map consisting of the most major numResults results. Returns the null AffineMap if the numResults is equal to zero. Returns the affineMap if numResults is greater or equals to number of results of the given affine map.
Returns the affine map consisting of the most minor numResults results. Returns the null AffineMap if the numResults is equal to zero. Returns the affineMap if numResults is greater or equals to number of results of the given affine map.
Checks whether the given affine map is an identity affine map. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an identity affine map on the most minor dimensions in the context. The affine map is owned by the context. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an affine map with a permutation expression and its size in the context. The permutation expression is a non-empty vector of integers. The elements of the permutation vector must be continuous from 0 and cannot be repeated (i.e. [1,2,0] is a valid permutation. [2,0] or [1,1,2] is an invalid permutation.) The affine map is owned by the context.
Prints an affine map by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Apply AffineExpr::replace(map) to each of the results and return a new new AffineMap with the new results and the specified number of dims and symbols.
Creates an instance of AnyQuantizedType with the given parameters in the same context as storageType and returns it. The instance is owned by the context.
Prints an attribute by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Takes an operation owned by the caller and inserts it as pos to the block. This is an expensive operation that scans the block linearly, prefer insertBefore/After instead.
Takes an operation owned by the caller and inserts it after the (non-owned) reference operation in the given block. If the reference is null, prepends the operation. Otherwise, the reference must belong to the block.
Takes an operation owned by the caller and inserts it before the (non-owned) reference operation in the given block. If the reference is null, appends the operation. Otherwise, the reference must belong to the block.
Prints a block by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates an instance of CalibratedQuantizedType with the given parameters in the same context as expressedType and returns it. The instance is owned by the context.
Attaches the diagnostic handler to the context. Handlers are invoked in the reverse order of attachment until one of them processes the diagnostic completely. When a handler is invoked it is passed the userData that was provided when it was attached. If non-NULL, deleteUserData is called once the system no longer needs to call the handler (for instance after the handler is detached or the context is destroyed). Returns an identifier that can be used to detach the handler.
Gets the dialect instance owned by the given context using the dialect namespace to identify it, loads (i.e., constructs the instance of) the dialect if necessary. If the dialect is not registered with the context, returns null. Use mlirContextLoad<Name>Dialect to load an unregistered dialect.
Returns whether the given fully-qualified operation (i.e. 'dialect.operation') is registered with the context. This will return true if the dialect is loaded and the operation is registered within the dialect.
Sets the thread pool of the context explicitly, enabling multithreading in the process. This API should be used to avoid re-creating thread pools in long-running applications that perform multiple compilations, see the C++ documentation for MLIRContext for details.
Creates an external MlirPass that calls the supplied callbacks using the supplied userData. If opName is empty, the pass is a generic operation pass. Otherwise it is an operation pass specific to the specified pass name.
Creates a dense elements attribute with the given shaped type from elements of a specific type. Expects the element type of the shaped type to match the data element type.
Creates a dense elements attribute with the given Shaped type and elements populated from a packed, row-major opaque buffer of contents.
The format of the raw buffer is a densely packed array of values that can be bitcast to the storage format of the element type specified. Types that are not byte aligned will be: - For bitwidth > 1: Rounded up to the next byte. - For bitwidth = 1: Packed into 8bit bytes with bits corresponding to the linear order of the shape type from MSB to LSB, padded to on the right.
A raw buffer of a single element (or for 1-bit, a byte of value 0 or 255) will be interpreted as a splat. User code should be prepared for additional, conformant patterns to be identified as splats in the future.
Creates a dense elements attribute that has the same data as the given dense elements attribute and a different shaped type. The new type must have the same total number of elements.
Gets the total number of elements in the given elements attribute. In order to iterate over the attribute, obtain its type, which must be a statically shaped type and use its sizes to build a multi-dimensional index.
Creates an ExecutionEngine for the provided ModuleOp. The ModuleOp is expected to be "translatable" to LLVM IR (only contains operations in dialects that implement the LLVMTranslationDialectInterface). The module ownership stays with the client and can be destroyed as soon as the call returns. optLevel is the optimization level to be used for transformation and code generation. LLVM passes at optLevel are run before code generation. The number and array of paths corresponding to shared libraries that will be loaded are specified via numPaths and sharedLibPaths respectively. TODO: figure out other options.
Invoke a native function in the execution engine by name with the arguments and result of the invoked function passed as an array of pointers. The function must have been tagged with the llvm.emit\_c\_interface attribute. Returns a failure if the execution fails for any reason (the function name can't be resolved for instance).
Infers the return types of the operation identified by its canonical given the arguments that will be supplied to its generic builder. Calls callback with the types of inferred arguments, potentially several times, on success. Returns failure otherwise.
Checks if two integer set objects are equal. This is a "shallow" comparison of two objects. Only the sets with some small number of constraints are uniqued and compare equal here. Set objects that represent the same integer set with different constraints may be considered non-equal by this check. Set difference followed by an (expensive) emptiness check should be used to check equivalence of the underlying integer sets.
Gets or creates a new integer set in the given context. The set is defined by a list of affine constraints, with the given number of input dimensions and symbols, which are treated as either equalities (eqFlags is 1) or inequalities (eqFlags is 0). Both constraints and eqFlags are expected to point to at least numConstraint consecutive values.
Prints an integer set by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Gets or creates a new integer set in which the values and dimensions of the given set are replaced with the given affine expressions. dimReplacements and symbolReplacements are expected to point to at least as many consecutive expressions as the given set has dimensions and symbols, respectively. The new set will have numResultDims and numResultSymbols dimensions and symbols, respectively.
Creates a LLVM DIDerivedType attribute. Note that dwarfAddressSpace is an optional field, where MLIR_CAPI_DWARF_ADDRESS_SPACE_NULL indicates null and non-negative values indicate a value present.
Creates an LLVM identified struct type with no body. If a struct type with this name already exists in the context, returns that type. Use mlirLLVMStructTypeIdentifiedNewGet to create a fresh struct type, potentially renaming it. The body should be set separatelty by calling mlirLLVMStructTypeSetBody, if it isn't set already.
Creates an LLVM identified struct type with no body and a name starting with the given prefix. If a struct with the exact name as the given prefix already exists, appends an unspecified suffix to the name so that the name is unique in context.
Creates an LLVM literal (unnamed) struct type. This may assert if the fields have types not compatible with the LLVM dialect. For a graceful failure, use the checked version.
Creates a name location owned by the given context. Providing null location for childLoc is allowed and if childLoc is null location, then the behavior is the same as having unknown child location.
Prints a location by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates a MemRef type with the given rank, shape, memory space and element type in the same context as the element type. The type has no affine maps, i.e. represents a default row-major contiguous memref. The type is owned by the context.
Creates a MemRef type with the given rank and shape, a potentially empty list of affine layout maps, the given memory space and element type, in the same context as element type. The type is owned by the context.
Returns the strides of the MemRef if the layout map is in strided form. Both strides and offset are out params. strides must point to pre-allocated memory of length equal to the rank of the memref.
Merge the symbols from other into target, potentially renaming them to avoid conflicts. Private symbols may be renamed during the merge, public symbols must have at most one declaration. A name conflict in public symbols is reported as an error before returning a failure.
Note that this clones the other operation unlike the C++ counterpart that takes ownership.
Add a pass and transfer ownership to the provided mlirOpPassManager. If the pass is not a generic operation pass or matching the type of the provided PassManager, a new OpPassManager is implicitly nested under the provided PassManager.
Parse a sequence of textual MLIR pass pipeline elements and add them to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Nest an OpPassManager under the provided OpPassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed.
Enables the elision of large elements attributes by printing a lexically valid but otherwise meaningless form instead of the element data. The largeElementLimit is used to configure what is considered to be a "large" ElementsAttr by providing an upper limit to the number of elements.
Enables the elision of large resources strings by omitting them from the dialect_resources section. The largeResourceLimit is used to configure what is considered to be a "large" resource by providing an upper limit to the string size.
Enable or disable printing of debug information (based on enable). If 'prettyForm' is set to true, debug information is printed in a more readable 'pretty' form. Note: The IR generated with 'prettyForm' is not parsable.
Use local scope when printing the operation. This allows for using the printer in a more localized and thread-safe setting, but may not necessarily be identical to what the IR will look like when dumping the full module.
Creates an opaque attribute in the given context associated with the dialect identified by its namespace. The attribute contains opaque byte data of the specified length (data need not be null-terminated).
Creates an opaque type in the given context associated with the dialect identified by its namespace. The type contains opaque byte data of the specified length (data need not be null-terminated).
Creates an operation and transfers ownership to the caller. Note that caller owned child objects are transferred in this call and must not be further used. Particularly, this applies to any regions added to the state (the implementation may invalidate any such pointers).
This call can fail under the following conditions, in which case, it will return a null operation and emit diagnostics: - Result type inference is enabled and cannot be performed.
Parses an operation, giving ownership to the caller. If parsing fails a null operation will be returned, and an error diagnostic emitted.
sourceStr may be either the text assembly format, or binary bytecode format. sourceName is used as the file name of the source; any IR without locations will get a FileLineColLoc location with sourceName as the file name.
Returns the number of attributes attached to the operation. Deprecated, please use mlirOperationGetNumInherentAttributes or mlirOperationGetNumDiscardableAttributes.
Returns true if this operation defines an inherent attribute with this name. Note: the attribute can be optional, so mlirOperationGetInherentAttributeByName can still return a null attribute.
Returns true if the operation identified by its canonical string name implements the interface identified by its TypeID in the given context. Note that interfaces may be attached to operations in some contexts and not others.
Moves the given operation immediately after the other operation in its parent block. The given operation may be owned by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Moves the given operation immediately before the other operation in its parent block. The given operation may be owner by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Prints an operation by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Removes an attribute by name. Returns false if the attribute was not found and true if removed. Deprecated, please use mlirOperationRemoveInherentAttributeByName or mlirOperationRemoveDiscardableAttributeByName.
Sets a discardable attribute by name, replacing the existing if it exists or adding a new one otherwise. The new attr Attribute is not allowed to be null, use mlirOperationRemoveDiscardableAttributeByName to remove an Attribute instead.
Enables result type inference for the operation under construction. If enabled, then the caller must not have called mlirOperationStateAddResults(). Note that if enabled, the mlirOperationCreate() call is failable: it will return a null operation on inference failure and will emit diagnostics.
Walks operation op in walkOrder and calls callback on that operation. *userData is passed to the callback as well and can be used to tunnel some context or other data into the callback.
Parse a textual MLIR pass pipeline and assign it to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Add a pass and transfer ownership to the provided top-level mlirPassManager. If the pass is not a generic operation pass or a ModulePass, a new OpPassManager is implicitly nested under the provided PassManager.
Nest an OpPassManager under the top-level PassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed. To further nest more OpPassManager under the newly returned one, see mlirOpPassManagerNest below.
Print a textual MLIR pass pipeline by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Casts from a type based on the expressed type of the given type to a corresponding type based on the given type. Returns a null type if the cast is not valid.
Casts from a type based on the storage type of the given type to a corresponding type based on the given type. Returns a null type if the cast is not valid.
Creates a tensor type of a fixed rank with the given shape, element type, and optional encoding in the same context as the element type. The type is owned by the context. Tensor types without any specific encoding field should assign mlirAttributeGetNull() to this parameter.
Takes a block owned by the caller and inserts it at pos to the given region. This is an expensive operation that linearly scans the region, prefer insertAfter/Before instead.
Takes a block owned by the caller and inserts it after the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, prepends the block to the region.
Takes a block owned by the caller and inserts it before the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, appends the block to the region.
Reset the insertion point to no location. Creating an operation without a set insertion point is an error, but this can still be useful when the current insertion point a builder refers to is being removed.
Add new block with 'argTypes' arguments and set the insertion point to the end of it. The block is placed before 'insertBefore'. locs contains the locations of the inserted arguments, and should match the size of argTypes.
This method is used to signal the end of an in-place modification of the given operation. This can only be called on operations that were provided to a call to startOpModification.
Inline the operations of block 'source' before the operation 'op'. The source block will be deleted and must have no uses. 'argValues' is used to replace the block arguments of 'source'
The source block must have no successors. Otherwise, the resulting IR would have unreachable operations.
Move the blocks that belong to "region" before the given position in another region "parent". The two regions must be different. The caller is responsible for creating or updating the operation transferring flow of control to the region and passing it the correct block arguments.
Inline the operations of block 'source' into the end of block 'dest'. The source block will be deleted and must have no uses. 'argValues' is used to replace the block arguments of 'source'
The dest block must have no successors. Otherwise, the resulting IR would have unreachable operation.
Find uses of from and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced) and that the from operation is about to be replaced.
Find uses of from and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced) and that the from operation is about to be replaced.
mlirRewriterBaseReplaceAllUsesExcept(rewriter, from, to, exceptedUser)
Find uses of from and replace them with to except if the user is exceptedUser. Also notify the listener about every in-place op modification (for every use that was replaced).
Find uses of from within block and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced). The optional allUsesReplaced flag is set to "true" if all uses were replaced.
Replace the results of the given (original) operation with the specified new op (replacement). The result types of the two ops must match. The original op is erased.
Replace the results of the given (original) operation with the specified list of values (replacements). The result types of the given op and the replacements must match. The original op is erased.
Sets the insertion point to the node after the specified value. If value has a defining operation, sets the insertion point to the node after such defining operation. This will cause subsequent insertions to go right after it. Otherwise, value is a BlockArgument. Sets the insertion point to the start of its block.
This method is used to notify the rewriter that an in-place operation modification is about to happen. A call to this function must be followed by a call to either finalizeOpModification or cancelOpModification. This is a minor efficiency win (it avoids creating a new operation and removing the old one) but also often allows simpler code in the client.
Sets the current debug type, similarly to -debug-only=type in the command-line tools. Note that global debug should be enabled for any output to be produced.
Sets multiple current debug types, similarly to `-debug-only=type1,type2" in the command-line tools. Note that global debug should be enabled for any output to be produced.
Returns the value indicating a dynamic stride or offset in a shaped type. Prefer mlirShapedTypeGetDynamicStrideOrOffset to direct comparisons with this value.
Creates a sparse elements attribute of the given shape from a list of indices and a list of associated values. Both lists are expected to be dense elements attributes with the same number of elements. The list of indices is expected to contain 64-bit integers. The attribute is created in the same context as the type.
Creates a symbol reference attribute in the given context referencing a symbol identified by the given string inside a list of nested references. Each of the references in the list must not be nested.
Inserts the given operation into the given symbol table. The operation must have the symbol trait. If the symbol table already has a symbol with the same name, renames the symbol being inserted to ensure name uniqueness. Note that this does not move the operation itself into the block of the symbol table operation, this should be done separately. Returns the name of the symbol after insertion.
Looks up a symbol with the given name in the given symbol table and returns the operation that corresponds to the symbol. If the symbol cannot be found, returns a null operation.
Attempt to replace all uses that are nested within the given operation of the given symbol 'oldSymbol' with the provided 'newSymbol'. This does not traverse into nested symbol tables. Will fail atomically if there are any unknown operations that may be potential symbol tables.
Walks all symbol table operations nested within, and including, op. For each symbol table operation, the provided callback is invoked with the op and a boolean signifying if the symbols within that symbol table can be treated as if all uses within the IR are visible to the caller. allSymUsesVisible identifies whether all of the symbol uses of symbols within op are visible.
Applies the transformation script starting at the given transform root operation to the given payload operation. The module containing the transform root as well as the transform options should be provided. The transform operation must implement TransformOpInterface and the module must be a ModuleOp. Returns the status of the application.
Translate operation that satisfies LLVM dialect module requirements into an LLVM IR module living in the given context. This translates operations from any dilalect that has a registered implementation of LLVMTranslationDialectInterface.
Returns
the generated LLVM IR Module from the translated MLIR module, it is owned by the caller.
Prints a location by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates an instance of UniformQuantizedPerAxisType with the given parameters in the same context as storageType and returns it. scales and zeroPoints point to nDims number of elements. The instance is owned by the context.
Creates an instance of UniformQuantizedType with the given parameters in the same context as storageType and returns it. The instance is owned by the context.
Unlike the typed accessors below, constructs the attribute with a raw data buffer and no type/alignment checking. Use a more strongly typed accessor if possible. If dataIsMutable is false, then an immutable AsmResourceBlob will be created and that passed data contents will be treated as const. If the deleter is non NULL, then it will be called when the data buffer can no longer be accessed (passing userData to it).
Prints a value by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Replace all uses of 'of' value with the 'with' value, updating anything in the IR that uses 'of' to use the other value instead. When this returns there are zero uses of 'of'.
Creates a vector type of the shape identified by its rank and dimensions, with the given element type in the same context as the element type. The type is owned by the context.
Creates a scalable vector type with the shape identified by its rank and dimensions. A subset of dimensions may be marked as scalable via the corresponding flag list, which is expected to have as many entries as the rank of the vector. The vector is created in the same context as the element type.
Given a MLIR module given as a string, calls the function identified by the func_name keyword parameter (default "main") with the provided arguments and return a tuple for each result of the call.
Within each process group in the process grid, concatenates the values of the operand tensor from each process along all_gather_dim and produces a result tensor.
Within each process group in the process grid, applies a reduction function computation to the values of the operand tensor from each process and produces a result tensor.
Within each process group in the process grid, splits the values of the operand tensor along split_dimension into parts, scatters the split parts between the processes, concatenates the scattered parts along concat_dimension and produces a result tensor.
Computes gradients of several inputs of BatchNormTrainingOp backpropagating from grad_output, and produces grad_operand, grad_scale and grad_offset tensors.
Computes mean and variance across batch and spatial dimensions and normalizes the operand tensor, for each feature in the feature_index dimension and produces output, batch_mean and batch_var tensors.
Performs a bitcast operation on operand tensor and produces a result tensor where the bits of the entire operand tensor are reinterpreted using the type of the result tensor.
Within each process group in the process grid, send the value of the operand tensor from the source process to the target processes and produce a result tensor.
Within each process group in the process grid, sends the value of the operand tensor from the source process to the target process and produces a result tensor.
Encapsulates an operation made up (composed) of other StableHLO operations, taking inputs and composite_attributes and producing results. The semantics of the op are implemented by the decomposition attribute. The composite op can be replaced with its decomposition without changing program semantics. In cases where inlining the decomposition does not provide the same op semantics, prefer using custom_call.
The version field (defaults to 0) is used to denote when a composite's semantics change.
This operation is functionally identical to broadcast_in_dim op, but the result shape is specified dynamically via output_dimensions.
It also accepts optional attributes to express static knowledge about the expanding behavior of dimensions. If not specified, all dimensions are assumed to be possibly expanding. The sets of dimensions that are known to be expanding and the set of dimensions that are known to be non-expanding must be disjoint and they must be a subset of the operand's dimensions.
Ensures that the operations that produce the operand are executed before any operations that depend on the result and prevents compiler transformations from moving operations across the barrier. Other than that, the operation is an identity, i.e. result = operand.
Performs element-wise conversion of operand to another floating-point type that uses exponent_bits and mantissa_bits and back to the original floating-point type and produces an output tensor.
Within each process group in the process grid, performs reduction, using computations, over the values of the operand tensor from each process, splits the reduction result along scatter_dimension into parts, and scatters the split parts between the processes to produce the result.
Returns an output filled with uniform random data and an updated output state output_state given an initial state initial_state using the pseudorandom number generator algorithm rng_algorithm.
Performs element-wise rounding towards the nearest integer, breaking ties towards the even integer, on the operand tensor and produces a result tensor.
Performs element-wise reciprocal square root operation on operand tensor and produces a result tensor, implementing the rSqrt operation from the IEEE-754 specification.
Produces results tensors which are equal to inputs tensors except that several slices specified by scatter_indices are updated with the values updates using update_computation.
Scatters the values from the source tensor using scatter based on the outcome of reduce_window of the input tensor using select and produces a result tensor.
%result = stablehlo.slice %operand [1:3, 4:8:2]
+ : (tensor<3x8xi64>) -> tensor<2x2xi64>
+
+// Same in generic form: the `1:3` above is mapped to the first entry in
+// `start_indices` and `limit_indices`, while `strides` is implicitly 1.
+// The `4:8:2` above is parsed into the second entry of `start_indices`,
+// `limit_indices` and `strides` respectively.
+%result = "stablehlo.slice" (%operand) {
+ start_indices = array<i64: 1, 4>,
+ limit_indices = array<i64: 3, 8>,
+ strides = array<i64: 1, 2>
+} : (tensor<3x8xi64>) -> tensor<2x2xi64>
Sorts a variadic number of tensors in inputs together, according to a custom comparator, along the given dimension and produces a variadic number of tensors as results.
The batch_dims attribute specifies the number of major batch dimensions (0 or more) that act like a multidimensional loop over both the operand and the index.
Performs element-wise conversion of quantized tensor operand to a floating-point tensor result according to the quantization parameters defined by the operand type.
Performs element-wise conversion of floating-point tensor or quantized tensor operand to a quantized tensor result according to the quantization parameters defined by the result type.
+
+
+
+
\ No newline at end of file
diff --git a/previews/PR363/assets/api_affine.md.5iiZk8SU.js b/previews/PR363/assets/api_affine.md.5iiZk8SU.js
new file mode 100644
index 00000000..8381284c
--- /dev/null
+++ b/previews/PR363/assets/api_affine.md.5iiZk8SU.js
@@ -0,0 +1,92 @@
+import{_ as l,c as o,j as a,a as s,G as t,a2 as i,B as p,o as r}from"./chunks/framework.2yyKLD8d.js";const A=JSON.parse('{"title":"Affine Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/affine.md","filePath":"api/affine.md","lastUpdated":null}'),d={name:"api/affine.md"},c={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"},h={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},v={class:"jldocstring custom-block"},R={class:"jldocstring custom-block"},x={class:"jldocstring custom-block"},I={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"};function j(M,e,L,w,T,V){const n=p("Badge");return r(),o("div",null,[e[41]||(e[41]=a("h1",{id:"Affine-Dialect",tabindex:"-1"},[s("Affine Dialect "),a("a",{class:"header-anchor",href:"#Affine-Dialect","aria-label":'Permalink to "Affine Dialect {#Affine-Dialect}"'},"")],-1)),e[42]||(e[42]=a("p",null,[s("Refer to the "),a("a",{href:"https://mlir.llvm.org/docs/Dialects/Affine/",target:"_blank",rel:"noreferrer"},"official documentation"),s(" for more details.")],-1)),a("details",c,[a("summary",null,[e[0]||(e[0]=a("a",{id:"Reactant.MLIR.Dialects.affine.apply-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.affine.apply-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.affine.apply")],-1)),e[1]||(e[1]=s()),t(n,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=i(`
apply
The affine.apply operation applies an affine mapping to a list of SSA values, yielding a single SSA value. The number of dimension and symbol arguments to affine.apply must be equal to the respective number of dimensional and symbolic inputs to the affine mapping; the affine mapping has to be one-dimensional, and so the affine.apply operation always returns one value. The input operands and result must all have ‘index’ type.
The affine.for operation represents an affine loop nest. It has one region containing its body. This region must contain one block that terminates with affine.yield. Note: when affine.for is printed in custom format, the terminator is omitted. The block has one argument of index type that represents the induction variable of the loop.
The affine.for operation executes its body a number of times iterating from a lower bound to an upper bound by a stride. The stride, represented by step, is a positive constant integer which defaults to "1" if not present. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound.
The lower and upper bounds of a affine.for operation are represented as an application of an affine mapping to a list of SSA values passed to the map. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols.
The affine mappings for the bounds may return multiple results, in which case the max/min keywords are required (for the lower/upper bound respectively), and the bound is the maximum/minimum of the returned values. There is no semantic ambiguity, but MLIR syntax requires the use of these keywords to make things more obvious to human readers.
Many upper and lower bounds are simple, so MLIR accepts two custom form syntaxes: the form that accepts a single 'ssa-id' (e.g. %N) is shorthand for applying that SSA value to a function that maps a single symbol to itself, e.g., ()[s]->(s)()[%N]. The integer literal form (e.g. -42) is shorthand for a nullary mapping function that returns the constant value (e.g. ()->(-42)()).
Example showing reverse iteration of the inner loop:
affine.for can also operate on loop-carried variables (iter_args) and return the final values after loop termination. The initial values of the variables are passed as additional SSA operands to the affine.for following the operands for the loop's lower and upper bounds. The operation's region has equivalent arguments for each variable representing the value of the variable at the current iteration.
The region must terminate with an affine.yield that passes all the current iteration variables to the next iteration, or to the affine.for's results if at the last iteration. For affine.for's that execute zero iterations, the 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.func @reduce(%buffer: memref<1024xf32>) -> (f32) {
+ // Initial sum set to 0.
+ %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
+ }
+ return %sum : f32
+}
If the affine.for defines any values, a yield terminator must be explicitly present. The number and types of the "affine.for" results must match the initial values in the iter_args binding and the yield operands.
The affine.if operation restricts execution to a subset of the loop iteration space defined by an integer set (a conjunction of affine constraints). A single affine.if may end with an optional else clause.
The condition of the affine.if is represented by an integer set (a conjunction of affine constraints), and the SSA values bound to the dimensions and symbols in the integer set. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols.
The affine.if operation contains two regions for the "then" and "else" clauses. affine.if may return results that are defined in its regions. The values defined are determined by which execution path is taken. Each region of the affine.if must contain a single block with no arguments, and be terminated by affine.yield. If affine.if defines no values, the affine.yield can be left out, and will be inserted implicitly. Otherwise, it must be explicit. If no values are defined, the else block may be empty (i.e. contain no blocks).
The affine.load op reads an element from a memref, where the index for each memref dimension is an affine expression of loop induction variables and symbols. The output of affine.load is a new value with the same type as the elements of the memref. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.min operation applies an affine mapping to a list of SSA values, and returns the minimum value of all result expressions. The number of dimension and symbol arguments to affine.min must be equal to the respective number of dimensional and symbolic inputs to the affine mapping; the affine.min operation always returns one value. The input operands and result must all have 'index' type.
The affine.parallel operation represents a hyper-rectangular affine parallel band, defining zero or more SSA values for its induction variables. It has one region capturing the parallel band body. The induction variables are represented as arguments of this region. These SSA values always have type index, which is the size of the machine word. The strides, represented by steps, are positive constant integers which defaults to "1" if not present. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound. The body region must contain exactly one block that terminates with affine.yield.
The lower and upper bounds of a parallel operation are represented as an application of an affine mapping to a list of SSA values passed to the map. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols. The list of expressions in each map is interpreted according to the respective bounds group attribute. If a single expression belongs to the group, then the result of this expression is taken as a lower(upper) bound of the corresponding loop induction variable. If multiple expressions belong to the group, then the lower(upper) bound is the max(min) of these values obtained from these expressions. The loop band has as many loops as elements in the group bounds attributes.
Each value yielded by affine.yield will be accumulated/reduced via one of the reduction methods defined in the AtomicRMWKind enum. The order of reduction is unspecified, and lowering may produce any valid ordering. Loops with a 0 trip count will produce as a result the identity value associated with each reduction (i.e. 0.0 for addf, 1.0 for mulf). Assign reductions for loops with a trip count != 1 produces undefined results.
Note: Calling AffineParallelOp::build will create the required region and block, and insert the required terminator if it is trivial (i.e. no values are yielded). Parsing will also create the required region, block, and terminator, even when they are missing from the textual representation.
The affine.prefetch op prefetches data from a memref location described with an affine subscript similar to affine.load, and has three attributes: a read/write specifier, a locality hint, and a cache type specifier as shown below:
mlir
affine.prefetch %0[%i, %j + 5], read, locality<3>, data : memref<400x400xi32>
The read/write specifier is either 'read' or 'write', the locality hint specifier ranges from locality<0> (no locality) to locality<3> (extremely local keep in cache). The cache type specifier is either 'data' or 'instr' and specifies whether the prefetch is performed on data cache or on instruction cache.
The affine.store op writes an element to a memref, where the index for each memref dimension is an affine expression of loop induction variables and symbols. The affine.store op stores a new value which is the same type as the elements of the memref. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.vector_load is the vector counterpart of affine.load. It reads a slice from a MemRef, supplied as its first operand, into a vector of the same base elemental type. The index for each memref dimension is an affine expression of loop induction variables and symbols. These indices determine the start position of the read within the memref. The shape of the return vector type determines the shape of the slice read from the memref. This slice is contiguous along the respective dimensions of the shape. Strided vector loads will be supported in the future. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.vector_store is the vector counterpart of affine.store. It writes a vector, supplied as its first operand, into a slice within a MemRef of the same base elemental type, supplied as its second operand. The index for each memref dimension is an affine expression of loop induction variables and symbols. These indices determine the start position of the write within the memref. The shape of th input vector determines the shape of the slice written to the memref. This slice is contiguous along the respective dimensions of the shape. Strided vector stores will be supported in the future. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
',12))]),a("details",k,[a("summary",null,[e[36]||(e[36]=a("a",{id:"Reactant.MLIR.Dialects.affine.yield-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.affine.yield-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.affine.yield")],-1)),e[37]||(e[37]=s()),t(n,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[38]||(e[38]=a("p",null,[a("code",null,"yield")],-1)),e[39]||(e[39]=a("p",null,[s("The "),a("code",null,"affine.yield"),s(" yields zero or more SSA values from an affine op region and terminates the region. The semantics of how the values yielded are used is defined by the parent operation. If "),a("code",null,"affine.yield"),s(" has any operands, the operands must match the parent operation's results. If the parent operation defines no values, then the "),a("code",null,"affine.yield"),s(" may be left out in the custom syntax and the builders will insert one implicitly. Otherwise, it has to be present in the syntax to indicate which values are yielded.")],-1)),e[40]||(e[40]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Affine.jl#L808-L820",target:"_blank",rel:"noreferrer"},"source")],-1))])])}const C=l(d,[["render",j]]);export{A as __pageData,C as default};
diff --git a/previews/PR363/assets/api_affine.md.5iiZk8SU.lean.js b/previews/PR363/assets/api_affine.md.5iiZk8SU.lean.js
new file mode 100644
index 00000000..8381284c
--- /dev/null
+++ b/previews/PR363/assets/api_affine.md.5iiZk8SU.lean.js
@@ -0,0 +1,92 @@
+import{_ as l,c as o,j as a,a as s,G as t,a2 as i,B as p,o as r}from"./chunks/framework.2yyKLD8d.js";const A=JSON.parse('{"title":"Affine Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/affine.md","filePath":"api/affine.md","lastUpdated":null}'),d={name:"api/affine.md"},c={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"},h={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},v={class:"jldocstring custom-block"},R={class:"jldocstring custom-block"},x={class:"jldocstring custom-block"},I={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"};function j(M,e,L,w,T,V){const n=p("Badge");return r(),o("div",null,[e[41]||(e[41]=a("h1",{id:"Affine-Dialect",tabindex:"-1"},[s("Affine Dialect "),a("a",{class:"header-anchor",href:"#Affine-Dialect","aria-label":'Permalink to "Affine Dialect {#Affine-Dialect}"'},"")],-1)),e[42]||(e[42]=a("p",null,[s("Refer to the "),a("a",{href:"https://mlir.llvm.org/docs/Dialects/Affine/",target:"_blank",rel:"noreferrer"},"official documentation"),s(" for more details.")],-1)),a("details",c,[a("summary",null,[e[0]||(e[0]=a("a",{id:"Reactant.MLIR.Dialects.affine.apply-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.affine.apply-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.affine.apply")],-1)),e[1]||(e[1]=s()),t(n,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=i(`
apply
The affine.apply operation applies an affine mapping to a list of SSA values, yielding a single SSA value. The number of dimension and symbol arguments to affine.apply must be equal to the respective number of dimensional and symbolic inputs to the affine mapping; the affine mapping has to be one-dimensional, and so the affine.apply operation always returns one value. The input operands and result must all have ‘index’ type.
The affine.for operation represents an affine loop nest. It has one region containing its body. This region must contain one block that terminates with affine.yield. Note: when affine.for is printed in custom format, the terminator is omitted. The block has one argument of index type that represents the induction variable of the loop.
The affine.for operation executes its body a number of times iterating from a lower bound to an upper bound by a stride. The stride, represented by step, is a positive constant integer which defaults to "1" if not present. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound.
The lower and upper bounds of a affine.for operation are represented as an application of an affine mapping to a list of SSA values passed to the map. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols.
The affine mappings for the bounds may return multiple results, in which case the max/min keywords are required (for the lower/upper bound respectively), and the bound is the maximum/minimum of the returned values. There is no semantic ambiguity, but MLIR syntax requires the use of these keywords to make things more obvious to human readers.
Many upper and lower bounds are simple, so MLIR accepts two custom form syntaxes: the form that accepts a single 'ssa-id' (e.g. %N) is shorthand for applying that SSA value to a function that maps a single symbol to itself, e.g., ()[s]->(s)()[%N]. The integer literal form (e.g. -42) is shorthand for a nullary mapping function that returns the constant value (e.g. ()->(-42)()).
Example showing reverse iteration of the inner loop:
affine.for can also operate on loop-carried variables (iter_args) and return the final values after loop termination. The initial values of the variables are passed as additional SSA operands to the affine.for following the operands for the loop's lower and upper bounds. The operation's region has equivalent arguments for each variable representing the value of the variable at the current iteration.
The region must terminate with an affine.yield that passes all the current iteration variables to the next iteration, or to the affine.for's results if at the last iteration. For affine.for's that execute zero iterations, the 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.func @reduce(%buffer: memref<1024xf32>) -> (f32) {
+ // Initial sum set to 0.
+ %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
+ }
+ return %sum : f32
+}
If the affine.for defines any values, a yield terminator must be explicitly present. The number and types of the "affine.for" results must match the initial values in the iter_args binding and the yield operands.
The affine.if operation restricts execution to a subset of the loop iteration space defined by an integer set (a conjunction of affine constraints). A single affine.if may end with an optional else clause.
The condition of the affine.if is represented by an integer set (a conjunction of affine constraints), and the SSA values bound to the dimensions and symbols in the integer set. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols.
The affine.if operation contains two regions for the "then" and "else" clauses. affine.if may return results that are defined in its regions. The values defined are determined by which execution path is taken. Each region of the affine.if must contain a single block with no arguments, and be terminated by affine.yield. If affine.if defines no values, the affine.yield can be left out, and will be inserted implicitly. Otherwise, it must be explicit. If no values are defined, the else block may be empty (i.e. contain no blocks).
The affine.load op reads an element from a memref, where the index for each memref dimension is an affine expression of loop induction variables and symbols. The output of affine.load is a new value with the same type as the elements of the memref. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.min operation applies an affine mapping to a list of SSA values, and returns the minimum value of all result expressions. The number of dimension and symbol arguments to affine.min must be equal to the respective number of dimensional and symbolic inputs to the affine mapping; the affine.min operation always returns one value. The input operands and result must all have 'index' type.
The affine.parallel operation represents a hyper-rectangular affine parallel band, defining zero or more SSA values for its induction variables. It has one region capturing the parallel band body. The induction variables are represented as arguments of this region. These SSA values always have type index, which is the size of the machine word. The strides, represented by steps, are positive constant integers which defaults to "1" if not present. The lower and upper bounds specify a half-open range: the range includes the lower bound but does not include the upper bound. The body region must contain exactly one block that terminates with affine.yield.
The lower and upper bounds of a parallel operation are represented as an application of an affine mapping to a list of SSA values passed to the map. The same restrictions hold for these SSA values as for all bindings of SSA values to dimensions and symbols. The list of expressions in each map is interpreted according to the respective bounds group attribute. If a single expression belongs to the group, then the result of this expression is taken as a lower(upper) bound of the corresponding loop induction variable. If multiple expressions belong to the group, then the lower(upper) bound is the max(min) of these values obtained from these expressions. The loop band has as many loops as elements in the group bounds attributes.
Each value yielded by affine.yield will be accumulated/reduced via one of the reduction methods defined in the AtomicRMWKind enum. The order of reduction is unspecified, and lowering may produce any valid ordering. Loops with a 0 trip count will produce as a result the identity value associated with each reduction (i.e. 0.0 for addf, 1.0 for mulf). Assign reductions for loops with a trip count != 1 produces undefined results.
Note: Calling AffineParallelOp::build will create the required region and block, and insert the required terminator if it is trivial (i.e. no values are yielded). Parsing will also create the required region, block, and terminator, even when they are missing from the textual representation.
The affine.prefetch op prefetches data from a memref location described with an affine subscript similar to affine.load, and has three attributes: a read/write specifier, a locality hint, and a cache type specifier as shown below:
mlir
affine.prefetch %0[%i, %j + 5], read, locality<3>, data : memref<400x400xi32>
The read/write specifier is either 'read' or 'write', the locality hint specifier ranges from locality<0> (no locality) to locality<3> (extremely local keep in cache). The cache type specifier is either 'data' or 'instr' and specifies whether the prefetch is performed on data cache or on instruction cache.
The affine.store op writes an element to a memref, where the index for each memref dimension is an affine expression of loop induction variables and symbols. The affine.store op stores a new value which is the same type as the elements of the memref. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.vector_load is the vector counterpart of affine.load. It reads a slice from a MemRef, supplied as its first operand, into a vector of the same base elemental type. The index for each memref dimension is an affine expression of loop induction variables and symbols. These indices determine the start position of the read within the memref. The shape of the return vector type determines the shape of the slice read from the memref. This slice is contiguous along the respective dimensions of the shape. Strided vector loads will be supported in the future. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
The affine.vector_store is the vector counterpart of affine.store. It writes a vector, supplied as its first operand, into a slice within a MemRef of the same base elemental type, supplied as its second operand. The index for each memref dimension is an affine expression of loop induction variables and symbols. These indices determine the start position of the write within the memref. The shape of th input vector determines the shape of the slice written to the memref. This slice is contiguous along the respective dimensions of the shape. Strided vector stores will be supported in the future. An affine expression of loop IVs and symbols must be specified for each dimension of the memref. The keyword symbol can be used to indicate SSA identifiers which are symbolic.
',12))]),a("details",k,[a("summary",null,[e[36]||(e[36]=a("a",{id:"Reactant.MLIR.Dialects.affine.yield-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.affine.yield-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.affine.yield")],-1)),e[37]||(e[37]=s()),t(n,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[38]||(e[38]=a("p",null,[a("code",null,"yield")],-1)),e[39]||(e[39]=a("p",null,[s("The "),a("code",null,"affine.yield"),s(" yields zero or more SSA values from an affine op region and terminates the region. The semantics of how the values yielded are used is defined by the parent operation. If "),a("code",null,"affine.yield"),s(" has any operands, the operands must match the parent operation's results. If the parent operation defines no values, then the "),a("code",null,"affine.yield"),s(" may be left out in the custom syntax and the builders will insert one implicitly. Otherwise, it has to be present in the syntax to indicate which values are yielded.")],-1)),e[40]||(e[40]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Affine.jl#L808-L820",target:"_blank",rel:"noreferrer"},"source")],-1))])])}const C=l(d,[["render",j]]);export{A as __pageData,C as default};
diff --git a/previews/PR363/assets/api_api.md.CxfFeNTq.js b/previews/PR363/assets/api_api.md.CxfFeNTq.js
new file mode 100644
index 00000000..f851d2bf
--- /dev/null
+++ b/previews/PR363/assets/api_api.md.CxfFeNTq.js
@@ -0,0 +1,27 @@
+import{_ as l,c as p,j as i,a,G as t,a2 as n,B as h,o}from"./chunks/framework.2yyKLD8d.js";const j=JSON.parse('{"title":"Core Reactant API","description":"","frontmatter":{},"headers":[],"relativePath":"api/api.md","filePath":"api/api.md","lastUpdated":null}'),r={name:"api/api.md"},d={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"},c={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"};function E(b,s,f,C,F,A){const e=h("Badge");return o(),p("div",null,[s[21]||(s[21]=i("h1",{id:"Core-Reactant-API",tabindex:"-1"},[a("Core Reactant API "),i("a",{class:"header-anchor",href:"#Core-Reactant-API","aria-label":'Permalink to "Core Reactant API {#Core-Reactant-API}"'},"")],-1)),s[22]||(s[22]=i("h2",{id:"Compile-API",tabindex:"-1"},[a("Compile API "),i("a",{class:"header-anchor",href:"#Compile-API","aria-label":'Permalink to "Compile API {#Compile-API}"'},"")],-1)),i("details",d,[i("summary",null,[s[0]||(s[0]=i("a",{id:"Reactant.Compiler.@compile",href:"#Reactant.Compiler.@compile"},[i("span",{class:"jlbinding"},"Reactant.Compiler.@compile")],-1)),s[1]||(s[1]=a()),t(e,{type:"info",class:"jlObjectType jlMacro",text:"Macro"})]),s[2]||(s[2]=n('
`,2))]),s[23]||(s[23]=i("h2",{id:"ReactantCore-API",tabindex:"-1"},[a("ReactantCore API "),i("a",{class:"header-anchor",href:"#ReactantCore-API","aria-label":'Permalink to "ReactantCore API {#ReactantCore-API}"'},"")],-1)),i("details",c,[i("summary",null,[s[6]||(s[6]=i("a",{id:"ReactantCore.@trace",href:"#ReactantCore.@trace"},[i("span",{class:"jlbinding"},"ReactantCore.@trace")],-1)),s[7]||(s[7]=a()),t(e,{type:"info",class:"jlObjectType jlMacro",text:"Macro"})]),s[8]||(s[8]=n(`
julia
@trace <expr>
Converts certain expressions like control flow into a Reactant friendly form. Importantly, if no traced value is found inside the expression, then there is no overhead.
Currently Supported
if conditions (with elseif and other niceties) (@trace if ...)
if statements with a preceeding assignment (@trace a = if ...) (note the positioning of the macro needs to be before the assignment and not before the if)
for statements with a single induction variable iterating over a syntactic StepRange of integers.
Special Considerations
Apply @trace only at the outermost if. Nested if statements will be automatically expanded into the correct form.
Extended Help
Caveats (Deviations from Core Julia Semantics)
New variables introduced
julia
@trace if x > 0
+ y = x + 1
+ p = 1
+else
+ y = x - 1
+end
In the outer scope p is not defined if x ≤ 0. However, for the traced version, it is defined and set to a dummy value.
Short Circuiting Operations
julia
@trace if x > 0 && z > 0
+ y = x + 1
+else
+ y = x - 1
+end
&& and || are short circuiting operations. In the traced version, we replace them with & and | respectively.
Type-Unstable Branches
julia
@trace if x > 0
+ y = 1.0f0
+else
+ y = 1.0
+end
This will not compile since y is a Float32 in one branch and a Float64 in the other. You need to ensure that all branches have the same type.
Another example is the following for loop which changes the type of x between iterations.
julia
x = ... # ConcreteRArray{Int64, 1}
+for i in 1f0:0.5f0:10f0
+ x = x .+ i # ConcreteRArray{Float32, 1}
+end
Certain Symbols are Reserved
Symbols like [😦😃, :nothing, :missing, :Inf, :Inf16, :Inf32, :Inf64, :Base, :Core] are not allowed as variables in @trace expressions. While certain cases might work but these are not guaranteed to work. For example, the following will not work:
julia
function fn(x)
+ nothing = sum(x)
+ @trace if nothing > 0
+ y = 1.0
+ else
+ y = 2.0
+ end
+ return y, nothing
+end
',2))]),s[25]||(s[25]=i("br",null,null,-1)),s[26]||(s[26]=i("h1",{id:"Internal-Functionality",tabindex:"-1"},[a("Internal Functionality "),i("a",{class:"header-anchor",href:"#Internal-Functionality","aria-label":'Permalink to "Internal Functionality {#Internal-Functionality}"'},"")],-1)),s[27]||(s[27]=i("div",{class:"danger custom-block"},[i("p",{class:"custom-block-title"},"Private"),i("p",null,"These functions are not part of the public API and are subject to change at any time.")],-1)),i("details",y,[i("summary",null,[s[12]||(s[12]=i("a",{id:"Reactant.Compiler.codegen_unflatten!",href:"#Reactant.Compiler.codegen_unflatten!"},[i("span",{class:"jlbinding"},"Reactant.Compiler.codegen_unflatten!")],-1)),s[13]||(s[13]=a()),t(e,{type:"info",class:"jlObjectType jlFunction",text:"Function"})]),s[14]||(s[14]=n('
julia
codegen_unflatten!
Generate Julia code to wrap the XLA buffers back into the output result datatypes. The name is due to its similarity to the unflatten function in jax.tree_util.register_pytree_node.
Generate Julia code to extract the XLA buffers from input arguments. The name is due to its similarity to the flatten function in jax.tree_util.register_pytree_node.
Arguments
linear_args: A list of arguments to be flattened.
Returns
flatten_names: A list of Symbols representing the names of the flattened arguments.
flatten_code: A list of Exprs to extract the XLA buffers from the input arguments.
Note
The linearized arguments do not directly refer to the are the arguments that have been flattened into a single list.
',5))])])}const D=l(r,[["render",E]]);export{j as __pageData,D as default};
diff --git a/previews/PR363/assets/api_api.md.CxfFeNTq.lean.js b/previews/PR363/assets/api_api.md.CxfFeNTq.lean.js
new file mode 100644
index 00000000..f851d2bf
--- /dev/null
+++ b/previews/PR363/assets/api_api.md.CxfFeNTq.lean.js
@@ -0,0 +1,27 @@
+import{_ as l,c as p,j as i,a,G as t,a2 as n,B as h,o}from"./chunks/framework.2yyKLD8d.js";const j=JSON.parse('{"title":"Core Reactant API","description":"","frontmatter":{},"headers":[],"relativePath":"api/api.md","filePath":"api/api.md","lastUpdated":null}'),r={name:"api/api.md"},d={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"},c={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"};function E(b,s,f,C,F,A){const e=h("Badge");return o(),p("div",null,[s[21]||(s[21]=i("h1",{id:"Core-Reactant-API",tabindex:"-1"},[a("Core Reactant API "),i("a",{class:"header-anchor",href:"#Core-Reactant-API","aria-label":'Permalink to "Core Reactant API {#Core-Reactant-API}"'},"")],-1)),s[22]||(s[22]=i("h2",{id:"Compile-API",tabindex:"-1"},[a("Compile API "),i("a",{class:"header-anchor",href:"#Compile-API","aria-label":'Permalink to "Compile API {#Compile-API}"'},"")],-1)),i("details",d,[i("summary",null,[s[0]||(s[0]=i("a",{id:"Reactant.Compiler.@compile",href:"#Reactant.Compiler.@compile"},[i("span",{class:"jlbinding"},"Reactant.Compiler.@compile")],-1)),s[1]||(s[1]=a()),t(e,{type:"info",class:"jlObjectType jlMacro",text:"Macro"})]),s[2]||(s[2]=n('
`,2))]),s[23]||(s[23]=i("h2",{id:"ReactantCore-API",tabindex:"-1"},[a("ReactantCore API "),i("a",{class:"header-anchor",href:"#ReactantCore-API","aria-label":'Permalink to "ReactantCore API {#ReactantCore-API}"'},"")],-1)),i("details",c,[i("summary",null,[s[6]||(s[6]=i("a",{id:"ReactantCore.@trace",href:"#ReactantCore.@trace"},[i("span",{class:"jlbinding"},"ReactantCore.@trace")],-1)),s[7]||(s[7]=a()),t(e,{type:"info",class:"jlObjectType jlMacro",text:"Macro"})]),s[8]||(s[8]=n(`
julia
@trace <expr>
Converts certain expressions like control flow into a Reactant friendly form. Importantly, if no traced value is found inside the expression, then there is no overhead.
Currently Supported
if conditions (with elseif and other niceties) (@trace if ...)
if statements with a preceeding assignment (@trace a = if ...) (note the positioning of the macro needs to be before the assignment and not before the if)
for statements with a single induction variable iterating over a syntactic StepRange of integers.
Special Considerations
Apply @trace only at the outermost if. Nested if statements will be automatically expanded into the correct form.
Extended Help
Caveats (Deviations from Core Julia Semantics)
New variables introduced
julia
@trace if x > 0
+ y = x + 1
+ p = 1
+else
+ y = x - 1
+end
In the outer scope p is not defined if x ≤ 0. However, for the traced version, it is defined and set to a dummy value.
Short Circuiting Operations
julia
@trace if x > 0 && z > 0
+ y = x + 1
+else
+ y = x - 1
+end
&& and || are short circuiting operations. In the traced version, we replace them with & and | respectively.
Type-Unstable Branches
julia
@trace if x > 0
+ y = 1.0f0
+else
+ y = 1.0
+end
This will not compile since y is a Float32 in one branch and a Float64 in the other. You need to ensure that all branches have the same type.
Another example is the following for loop which changes the type of x between iterations.
julia
x = ... # ConcreteRArray{Int64, 1}
+for i in 1f0:0.5f0:10f0
+ x = x .+ i # ConcreteRArray{Float32, 1}
+end
Certain Symbols are Reserved
Symbols like [😦😃, :nothing, :missing, :Inf, :Inf16, :Inf32, :Inf64, :Base, :Core] are not allowed as variables in @trace expressions. While certain cases might work but these are not guaranteed to work. For example, the following will not work:
julia
function fn(x)
+ nothing = sum(x)
+ @trace if nothing > 0
+ y = 1.0
+ else
+ y = 2.0
+ end
+ return y, nothing
+end
',2))]),s[25]||(s[25]=i("br",null,null,-1)),s[26]||(s[26]=i("h1",{id:"Internal-Functionality",tabindex:"-1"},[a("Internal Functionality "),i("a",{class:"header-anchor",href:"#Internal-Functionality","aria-label":'Permalink to "Internal Functionality {#Internal-Functionality}"'},"")],-1)),s[27]||(s[27]=i("div",{class:"danger custom-block"},[i("p",{class:"custom-block-title"},"Private"),i("p",null,"These functions are not part of the public API and are subject to change at any time.")],-1)),i("details",y,[i("summary",null,[s[12]||(s[12]=i("a",{id:"Reactant.Compiler.codegen_unflatten!",href:"#Reactant.Compiler.codegen_unflatten!"},[i("span",{class:"jlbinding"},"Reactant.Compiler.codegen_unflatten!")],-1)),s[13]||(s[13]=a()),t(e,{type:"info",class:"jlObjectType jlFunction",text:"Function"})]),s[14]||(s[14]=n('
julia
codegen_unflatten!
Generate Julia code to wrap the XLA buffers back into the output result datatypes. The name is due to its similarity to the unflatten function in jax.tree_util.register_pytree_node.
Generate Julia code to extract the XLA buffers from input arguments. The name is due to its similarity to the flatten function in jax.tree_util.register_pytree_node.
Arguments
linear_args: A list of arguments to be flattened.
Returns
flatten_names: A list of Symbols representing the names of the flattened arguments.
flatten_code: A list of Exprs to extract the XLA buffers from the input arguments.
Note
The linearized arguments do not directly refer to the are the arguments that have been flattened into a single list.
',5))])])}const D=l(r,[["render",E]]);export{j as __pageData,D as default};
diff --git a/previews/PR363/assets/api_arith.md.DbkUJ4WC.js b/previews/PR363/assets/api_arith.md.DbkUJ4WC.js
new file mode 100644
index 00000000..8abe164c
--- /dev/null
+++ b/previews/PR363/assets/api_arith.md.DbkUJ4WC.js
@@ -0,0 +1,186 @@
+import{_ as l,c as o,j as a,a as t,G as n,a2 as i,B as p,o as r}from"./chunks/framework.2yyKLD8d.js";const pe=JSON.parse('{"title":"Arithmetic Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/arith.md","filePath":"api/arith.md","lastUpdated":null}'),c={name:"api/arith.md"},d={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"},h={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"},R={class:"jldocstring custom-block"},v={class:"jldocstring custom-block"},I={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},M={class:"jldocstring custom-block"},L={class:"jldocstring custom-block"},j={class:"jldocstring custom-block"},x={class:"jldocstring custom-block"},D={class:"jldocstring custom-block"},T={class:"jldocstring custom-block"},w={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"},V={class:"jldocstring custom-block"},A={class:"jldocstring custom-block"},E={class:"jldocstring custom-block"},q={class:"jldocstring custom-block"},z={class:"jldocstring custom-block"},C={class:"jldocstring custom-block"},S={class:"jldocstring custom-block"},O={class:"jldocstring custom-block"},N={class:"jldocstring custom-block"},W={class:"jldocstring custom-block"},B={class:"jldocstring custom-block"},P={class:"jldocstring custom-block"},U={class:"jldocstring custom-block"},F={class:"jldocstring custom-block"},G={class:"jldocstring custom-block"},$={class:"jldocstring custom-block"},J={class:"jldocstring custom-block"},H={class:"jldocstring custom-block"},K={class:"jldocstring custom-block"},Q={class:"jldocstring custom-block"},X={class:"jldocstring custom-block"},Y={class:"jldocstring custom-block"},Z={class:"jldocstring custom-block"},_={class:"jldocstring custom-block"},ee={class:"jldocstring custom-block"};function ae(te,e,se,ne,ie,le){const s=p("Badge");return r(),o("div",null,[e[153]||(e[153]=a("h1",{id:"Arithmetic-Dialect",tabindex:"-1"},[t("Arithmetic Dialect "),a("a",{class:"header-anchor",href:"#Arithmetic-Dialect","aria-label":'Permalink to "Arithmetic Dialect {#Arithmetic-Dialect}"'},"")],-1)),e[154]||(e[154]=a("p",null,[t("Refer to the "),a("a",{href:"https://mlir.llvm.org/docs/Dialects/ArithOps/",target:"_blank",rel:"noreferrer"},"official documentation"),t(" for more details.")],-1)),a("details",d,[a("summary",null,[e[0]||(e[0]=a("a",{id:"Reactant.MLIR.Dialects.arith.addf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.addf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.addf")],-1)),e[1]||(e[1]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=i(`
addf
The addf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit addition on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the addition modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The addi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Performs (N+1)-bit addition on zero-extended operands. Returns two results: the N-bit sum (same type as both operands), and the overflow bit (boolean-like), where 1 indicates unsigned addition overflow, while 0 indicates no overflow.
The andi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
`,5))]),a("details",g,[a("summary",null,[e[12]||(e[12]=a("a",{id:"Reactant.MLIR.Dialects.arith.bitcast-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.bitcast-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.bitcast")],-1)),e[13]||(e[13]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[14]||(e[14]=a("p",null,[a("code",null,"bitcast")],-1)),e[15]||(e[15]=a("p",null,"Bitcast an integer or floating point value to an integer or floating point value of equal bit width. When operating on vectors, casts elementwise.",-1)),e[16]||(e[16]=a("p",null,[t("Note that this implements a logical bitcast independent of target endianness. This allows constant folding without target information and is consitent with the bitcast constant folders in LLVM (see "),a("a",{href:"https://github.com/llvm/llvm-project/blob/18c19414eb/llvm/lib/IR/ConstantFold.cpp#L168",target:"_blank",rel:"noreferrer"},"https://github.com/llvm/llvm-project/blob/18c19414eb/llvm/lib/IR/ConstantFold.cpp#L168"),t(") For targets where the source and target type have the same endianness (which is the standard), this cast will also change no bits at runtime, but it may still require an operation, for example if the machine has different floating point and integer register files. For targets that have a different endianness for the source and target types (e.g. float is big-endian and integer is little-endian) a proper lowering would add operations to swap the order of words in addition to the bitcast.")],-1)),e[17]||(e[17]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L214-L231",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",f,[a("summary",null,[e[18]||(e[18]=a("a",{id:"Reactant.MLIR.Dialects.arith.ceildivsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.ceildivsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.ceildivsi")],-1)),e[19]||(e[19]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[20]||(e[20]=i(`
ceildivsi
Signed integer division. Rounds towards positive infinity, i.e. 7 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Unsigned integer division. Rounds towards positive infinity. Treats the leading bit as the most significant, i.e. for i16 given two's complement representation, 6 / -2 = 6 / (2^16 - 2) = 1.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
The cmpf operation compares its two operands according to the float comparison rules and the predicate specified by the respective attribute. The predicate defines the type of comparison: (un)orderedness, (in)equality and signed less/greater than (or equal to) as well as predicates that are always true or false. The operands must have the same type, and this type must be a float type, or a vector or tensor thereof. The result is an i1, or a vector/tensor thereof having the same shape as the inputs. Unlike cmpi, the operands are always treated as signed. The u prefix indicates unordered comparison, not unsigned comparison, so "une" means unordered or not equal. For the sake of readability by humans, custom assembly form for the operation uses a string-typed attribute for the predicate. The value of this attribute corresponds to lower-cased name of the predicate constant, e.g., "one" means "ordered not equal". The string representation of the attribute is merely a syntactic sugar and is converted to an integer attribute by the parser.
The cmpi operation is a generic comparison for integer-like types. Its two arguments can be integers, vectors or tensors thereof as long as their types match. The operation produces an i1 for the former case, a vector or a tensor of i1 with the same shape as inputs in the other cases.
Its first argument is an attribute that defines which type of comparison is performed. The following comparisons are supported:
equal (mnemonic: "eq"; integer value: 0)
not equal (mnemonic: "ne"; integer value: 1)
signed less than (mnemonic: "slt"; integer value: 2)
signed less than or equal (mnemonic: "sle"; integer value: 3)
signed greater than (mnemonic: "sgt"; integer value: 4)
signed greater than or equal (mnemonic: "sge"; integer value: 5)
unsigned less than (mnemonic: "ult"; integer value: 6)
unsigned less than or equal (mnemonic: "ule"; integer value: 7)
unsigned greater than (mnemonic: "ugt"; integer value: 8)
unsigned greater than or equal (mnemonic: "uge"; integer value: 9)
The result is 1 if the comparison is true and 0 otherwise. For vector or tensor operands, the comparison is performed elementwise and the element of the result indicates whether the comparison is true for the operand elements with the same indices as those of the result.
Note: while the custom assembly form uses strings, the actual underlying attribute has integer type (or rather enum class in C++ code) as seen from the generic assembly form. String literals are used to improve readability of the IR by humans.
This operation only applies to integer-like operands, but not floats. The main reason being that comparison operations have diverging sets of attributes: integers require sign specification while floats require various floating point-related particularities, e.g., -ffast-math behavior, IEEE754 compliance, etc (rationale). The type of comparison is specified as attribute to avoid introducing ten similar operations, taking into account that they are often implemented using the same operation downstream (rationale). The separation between signed and unsigned order comparisons is necessary because of integers being signless. The comparison operation must know how to interpret values with the foremost bit being set: negatives in two's complement or large positives (rationale).
Example
mlir
// Custom form of scalar "signed less than" comparison.
+%x = arith.cmpi slt, %lhs, %rhs : i32
+
+// Generic form of the same operation.
+%x = "arith.cmpi"(%lhs, %rhs) {predicate = 2 : i64} : (i32, i32) -> i1
+
+// Custom form of vector equality comparison.
+%x = arith.cmpi eq, %lhs, %rhs : vector<4xi64>
+
+// Generic form of the same operation.
+%x = "arith.cmpi"(%lhs, %rhs) {predicate = 0 : i64}
+ : (vector<4xi64>, vector<4xi64>) -> vector<4xi1>
The constant operation produces an SSA value equal to some integer or floating-point constant specified by an attribute. This is the way MLIR forms simple integer and floating point constants.
Signed integer division. Rounds towards zero. Treats the leading bit as sign, i.e. 6 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Unsigned integer division. Rounds towards zero. Treats the leading bit as the most significant, i.e. for i16 given two's complement representation, 6 / -2 = 6 / (2^16 - 2) = 0.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
`,6))]),a("details",L,[a("summary",null,[e[39]||(e[39]=a("a",{id:"Reactant.MLIR.Dialects.arith.extf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.extf-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.extf")],-1)),e[40]||(e[40]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[41]||(e[41]=a("p",null,[a("code",null,"extf")],-1)),e[42]||(e[42]=a("p",null,"Cast a floating-point value to a larger floating-point-typed value. The destination type must to be strictly wider than the source type. When operating on vectors, casts elementwise.",-1)),e[43]||(e[43]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L632-L638",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",j,[a("summary",null,[e[44]||(e[44]=a("a",{id:"Reactant.MLIR.Dialects.arith.extsi-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.extsi-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.extsi")],-1)),e[45]||(e[45]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[46]||(e[46]=i(`
extsi
The integer sign extension operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be larger than the input bit-width (N > M). The top-most (N - M) bits of the output are filled with copies of the most-significant bit of the input.
Example
mlir
%1 = arith.constant 5 : i3 // %1 is 0b101
+%2 = arith.extsi %1 : i3 to i6 // %2 is 0b111101
+%3 = arith.constant 2 : i3 // %3 is 0b010
+%4 = arith.extsi %3 : i3 to i6 // %4 is 0b000010
+
+%5 = arith.extsi %0 : vector<2 x i32> to vector<2 x i64>
The integer zero extension operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be larger than the input bit-width (N > M). The top-most (N - M) bits of the output are filled with zeros.
Example
mlir
%1 = arith.constant 5 : i3 // %1 is 0b101
+ %2 = arith.extui %1 : i3 to i6 // %2 is 0b000101
+ %3 = arith.constant 2 : i3 // %3 is 0b010
+ %4 = arith.extui %3 : i3 to i6 // %4 is 0b000010
+
+ %5 = arith.extui %0 : vector<2 x i32> to vector<2 x i64>
Signed integer division. Rounds towards negative infinity, i.e. 5 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
`,6))]),a("details",T,[a("summary",null,[e[53]||(e[53]=a("a",{id:"Reactant.MLIR.Dialects.arith.fptosi-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.fptosi-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.fptosi")],-1)),e[54]||(e[54]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[55]||(e[55]=a("p",null,[a("code",null,"fptosi")],-1)),e[56]||(e[56]=a("p",null,"Cast from a value interpreted as floating-point to the nearest (rounding towards zero) signed integer value. When operating on vectors, casts elementwise.",-1)),e[57]||(e[57]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L736-L742",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",w,[a("summary",null,[e[58]||(e[58]=a("a",{id:"Reactant.MLIR.Dialects.arith.fptoui-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.fptoui-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.fptoui")],-1)),e[59]||(e[59]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[60]||(e[60]=a("p",null,[a("code",null,"fptoui")],-1)),e[61]||(e[61]=a("p",null,"Cast from a value interpreted as floating-point to the nearest (rounding towards zero) unsigned integer value. When operating on vectors, casts elementwise.",-1)),e[62]||(e[62]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L762-L768",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",k,[a("summary",null,[e[63]||(e[63]=a("a",{id:"Reactant.MLIR.Dialects.arith.index_cast-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.index_cast-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.index_cast")],-1)),e[64]||(e[64]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[65]||(e[65]=a("p",null,[a("code",null,"index_cast")],-1)),e[66]||(e[66]=a("p",null,"Casts between scalar or vector integers and corresponding 'index' scalar or vectors. Index is an integer of platform-specific bit width. If casting to a wider integer, the value is sign-extended. If casting to a narrower integer, the value is truncated.",-1)),e[67]||(e[67]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L828-L835",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",V,[a("summary",null,[e[68]||(e[68]=a("a",{id:"Reactant.MLIR.Dialects.arith.index_castui-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.index_castui-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.index_castui")],-1)),e[69]||(e[69]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[70]||(e[70]=a("p",null,[a("code",null,"index_castui")],-1)),e[71]||(e[71]=a("p",null,"Casts between scalar or vector integers and corresponding 'index' scalar or vectors. Index is an integer of platform-specific bit width. If casting to a wider integer, the value is zero-extended. If casting to a narrower integer, the value is truncated.",-1)),e[72]||(e[72]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L855-L862",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",A,[a("summary",null,[e[73]||(e[73]=a("a",{id:"Reactant.MLIR.Dialects.arith.maximumf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.maximumf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.maximumf")],-1)),e[74]||(e[74]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[75]||(e[75]=i(`
maximumf
Returns the maximum of the two arguments, treating -0.0 as less than +0.0. If one of the arguments is NaN, then the result is also NaN.
Returns the maximum of the two arguments. If the arguments are -0.0 and +0.0, then the result is either of them. If one of the arguments is NaN, then the result is the other argument.
Returns the minimum of the two arguments. If the arguments are -0.0 and +0.0, then the result is either of them. If one of the arguments is NaN, then the result is the other argument.
The mulf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit multiplication on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the multiplication modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The muli operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Performs (2*N)-bit multiplication on sign-extended operands. Returns two N-bit results: the low and the high halves of the product. The low half has the same value as the result of regular multiplication arith.muli with the same operands.
Performs (2*N)-bit multiplication on zero-extended operands. Returns two N-bit results: the low and the high halves of the product. The low half has the same value as the result of regular multiplication arith.muli with the same operands.
The negf operation computes the negation of a given value. It takes one operand and returns one result of the same type. This type may be a float scalar type, a vector whose element type is float, or a tensor of floats. It has no standard attributes.
The ori operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
`,5))]),a("details",P,[a("summary",null,[e[103]||(e[103]=a("a",{id:"Reactant.MLIR.Dialects.arith.remf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.remf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.remf")],-1)),e[104]||(e[104]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[105]||(e[105]=a("p",null,[a("code",null,"remf")],-1)),e[106]||(e[106]=a("p",null,"Returns the floating point division remainder. The remainder has the same sign as the dividend (lhs operand).",-1)),e[107]||(e[107]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1431-L1436",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",U,[a("summary",null,[e[108]||(e[108]=a("a",{id:"Reactant.MLIR.Dialects.arith.remsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.remsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.remsi")],-1)),e[109]||(e[109]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[110]||(e[110]=i(`
remsi
Signed integer division remainder. Treats the leading bit as sign, i.e. 6 % -2 = 0.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
The arith.select operation chooses one value based on a binary condition supplied as its first operand.
If the value of the first operand (the condition) is 1, then the second operand is returned, and the third operand is ignored, even if it was poison.
If the value of the first operand (the condition) is 0, then the third operand is returned, and the second operand is ignored, even if it was poison.
If the value of the first operand (the condition) is poison, then the operation returns poison.
The operation applies to vectors and tensors elementwise given the shape of all operands is identical. The choice is made for each element individually based on the value at the same position as the element in the condition operand. If an i1 is provided as the condition, the entire vector or tensor is chosen.
Example
mlir
// Custom form of scalar selection.
+%x = arith.select %cond, %true, %false : i32
+
+// Generic form of the same operation.
+%x = "arith.select"(%cond, %true, %false) : (i1, i32, i32) -> i32
+
+// Element-wise vector selection.
+%vx = arith.select %vcond, %vtrue, %vfalse : vector<42xi1>, vector<42xf32>
+
+// Full vector selection.
+%vx = arith.select %cond, %vtrue, %vfalse : vector<42xf32>
The shli operation shifts the integer value of the first operand to the left by the integer value of the second operand. The second operand is interpreted as unsigned. The low order bits are 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.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
The shrsi operation shifts an integer value of the first operand to the right by the value of the second operand. The first operand is interpreted as signed, and the second operand is interpreter as unsigned. The high order bits in the output are filled with copies of the most-significant bit of the shifted value (which means that the sign of the value is preserved). If the value of the second operand is greater or equal than bitwidth of the first operand, then the operation returns poison.
The shrui operation shifts an integer value of the first operand to the right by the value of the second operand. The first operand is interpreted as unsigned, and the second operand is interpreted as unsigned. The high order bits are always 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.
`,5))]),a("details",K,[a("summary",null,[e[126]||(e[126]=a("a",{id:"Reactant.MLIR.Dialects.arith.sitofp-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.sitofp-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.sitofp")],-1)),e[127]||(e[127]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[128]||(e[128]=a("p",null,[a("code",null,"sitofp")],-1)),e[129]||(e[129]=a("p",null,"Cast from a value interpreted as a signed integer to the corresponding floating-point value. If the value cannot be exactly represented, it is rounded using the default rounding mode. When operating on vectors, casts elementwise.",-1)),e[130]||(e[130]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1554-L1561",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",Q,[a("summary",null,[e[131]||(e[131]=a("a",{id:"Reactant.MLIR.Dialects.arith.subf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.subf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.subf")],-1)),e[132]||(e[132]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[133]||(e[133]=i(`
subf
The subf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit subtraction on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the subtraction modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The subi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
`,7))]),a("details",Y,[a("summary",null,[e[137]||(e[137]=a("a",{id:"Reactant.MLIR.Dialects.arith.truncf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.truncf-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.truncf")],-1)),e[138]||(e[138]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[139]||(e[139]=a("p",null,[a("code",null,"truncf")],-1)),e[140]||(e[140]=a("p",null,"Truncate a floating-point value to a smaller floating-point-typed value. The destination type must be strictly narrower than the source type. If the value cannot be exactly represented, it is rounded using the provided rounding mode or the default one if no rounding mode is provided. When operating on vectors, casts elementwise.",-1)),e[141]||(e[141]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1827-L1835",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",Z,[a("summary",null,[e[142]||(e[142]=a("a",{id:"Reactant.MLIR.Dialects.arith.trunci-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.trunci-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.trunci")],-1)),e[143]||(e[143]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[144]||(e[144]=i(`
trunci
The integer truncation operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be smaller than the input bit-width (N < M). The top-most (N - M) bits of the input are discarded.
Example
mlir
%1 = arith.constant 21 : i5 // %1 is 0b10101
+ %2 = arith.trunci %1 : i5 to i4 // %2 is 0b0101
+ %3 = arith.trunci %1 : i5 to i3 // %3 is 0b101
+
+ %5 = arith.trunci %0 : vector<2 x i32> to vector<2 x i16>
`,5))]),a("details",_,[a("summary",null,[e[145]||(e[145]=a("a",{id:"Reactant.MLIR.Dialects.arith.uitofp-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.uitofp-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.uitofp")],-1)),e[146]||(e[146]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[147]||(e[147]=a("p",null,[a("code",null,"uitofp")],-1)),e[148]||(e[148]=a("p",null,"Cast from a value interpreted as unsigned integer to the corresponding floating-point value. If the value cannot be exactly represented, it is rounded using the default rounding mode. When operating on vectors, casts elementwise.",-1)),e[149]||(e[149]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1897-L1904",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",ee,[a("summary",null,[e[150]||(e[150]=a("a",{id:"Reactant.MLIR.Dialects.arith.xori-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.xori-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.xori")],-1)),e[151]||(e[151]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[152]||(e[152]=i(`
xori
The xori operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
`,5))])])}const re=l(c,[["render",ae]]);export{pe as __pageData,re as default};
diff --git a/previews/PR363/assets/api_arith.md.DbkUJ4WC.lean.js b/previews/PR363/assets/api_arith.md.DbkUJ4WC.lean.js
new file mode 100644
index 00000000..8abe164c
--- /dev/null
+++ b/previews/PR363/assets/api_arith.md.DbkUJ4WC.lean.js
@@ -0,0 +1,186 @@
+import{_ as l,c as o,j as a,a as t,G as n,a2 as i,B as p,o as r}from"./chunks/framework.2yyKLD8d.js";const pe=JSON.parse('{"title":"Arithmetic Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/arith.md","filePath":"api/arith.md","lastUpdated":null}'),c={name:"api/arith.md"},d={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"},h={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"},R={class:"jldocstring custom-block"},v={class:"jldocstring custom-block"},I={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},M={class:"jldocstring custom-block"},L={class:"jldocstring custom-block"},j={class:"jldocstring custom-block"},x={class:"jldocstring custom-block"},D={class:"jldocstring custom-block"},T={class:"jldocstring custom-block"},w={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"},V={class:"jldocstring custom-block"},A={class:"jldocstring custom-block"},E={class:"jldocstring custom-block"},q={class:"jldocstring custom-block"},z={class:"jldocstring custom-block"},C={class:"jldocstring custom-block"},S={class:"jldocstring custom-block"},O={class:"jldocstring custom-block"},N={class:"jldocstring custom-block"},W={class:"jldocstring custom-block"},B={class:"jldocstring custom-block"},P={class:"jldocstring custom-block"},U={class:"jldocstring custom-block"},F={class:"jldocstring custom-block"},G={class:"jldocstring custom-block"},$={class:"jldocstring custom-block"},J={class:"jldocstring custom-block"},H={class:"jldocstring custom-block"},K={class:"jldocstring custom-block"},Q={class:"jldocstring custom-block"},X={class:"jldocstring custom-block"},Y={class:"jldocstring custom-block"},Z={class:"jldocstring custom-block"},_={class:"jldocstring custom-block"},ee={class:"jldocstring custom-block"};function ae(te,e,se,ne,ie,le){const s=p("Badge");return r(),o("div",null,[e[153]||(e[153]=a("h1",{id:"Arithmetic-Dialect",tabindex:"-1"},[t("Arithmetic Dialect "),a("a",{class:"header-anchor",href:"#Arithmetic-Dialect","aria-label":'Permalink to "Arithmetic Dialect {#Arithmetic-Dialect}"'},"")],-1)),e[154]||(e[154]=a("p",null,[t("Refer to the "),a("a",{href:"https://mlir.llvm.org/docs/Dialects/ArithOps/",target:"_blank",rel:"noreferrer"},"official documentation"),t(" for more details.")],-1)),a("details",d,[a("summary",null,[e[0]||(e[0]=a("a",{id:"Reactant.MLIR.Dialects.arith.addf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.addf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.addf")],-1)),e[1]||(e[1]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=i(`
addf
The addf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit addition on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the addition modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The addi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Performs (N+1)-bit addition on zero-extended operands. Returns two results: the N-bit sum (same type as both operands), and the overflow bit (boolean-like), where 1 indicates unsigned addition overflow, while 0 indicates no overflow.
The andi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
`,5))]),a("details",g,[a("summary",null,[e[12]||(e[12]=a("a",{id:"Reactant.MLIR.Dialects.arith.bitcast-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.bitcast-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.bitcast")],-1)),e[13]||(e[13]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[14]||(e[14]=a("p",null,[a("code",null,"bitcast")],-1)),e[15]||(e[15]=a("p",null,"Bitcast an integer or floating point value to an integer or floating point value of equal bit width. When operating on vectors, casts elementwise.",-1)),e[16]||(e[16]=a("p",null,[t("Note that this implements a logical bitcast independent of target endianness. This allows constant folding without target information and is consitent with the bitcast constant folders in LLVM (see "),a("a",{href:"https://github.com/llvm/llvm-project/blob/18c19414eb/llvm/lib/IR/ConstantFold.cpp#L168",target:"_blank",rel:"noreferrer"},"https://github.com/llvm/llvm-project/blob/18c19414eb/llvm/lib/IR/ConstantFold.cpp#L168"),t(") For targets where the source and target type have the same endianness (which is the standard), this cast will also change no bits at runtime, but it may still require an operation, for example if the machine has different floating point and integer register files. For targets that have a different endianness for the source and target types (e.g. float is big-endian and integer is little-endian) a proper lowering would add operations to swap the order of words in addition to the bitcast.")],-1)),e[17]||(e[17]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L214-L231",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",f,[a("summary",null,[e[18]||(e[18]=a("a",{id:"Reactant.MLIR.Dialects.arith.ceildivsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.ceildivsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.ceildivsi")],-1)),e[19]||(e[19]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[20]||(e[20]=i(`
ceildivsi
Signed integer division. Rounds towards positive infinity, i.e. 7 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Unsigned integer division. Rounds towards positive infinity. Treats the leading bit as the most significant, i.e. for i16 given two's complement representation, 6 / -2 = 6 / (2^16 - 2) = 1.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
The cmpf operation compares its two operands according to the float comparison rules and the predicate specified by the respective attribute. The predicate defines the type of comparison: (un)orderedness, (in)equality and signed less/greater than (or equal to) as well as predicates that are always true or false. The operands must have the same type, and this type must be a float type, or a vector or tensor thereof. The result is an i1, or a vector/tensor thereof having the same shape as the inputs. Unlike cmpi, the operands are always treated as signed. The u prefix indicates unordered comparison, not unsigned comparison, so "une" means unordered or not equal. For the sake of readability by humans, custom assembly form for the operation uses a string-typed attribute for the predicate. The value of this attribute corresponds to lower-cased name of the predicate constant, e.g., "one" means "ordered not equal". The string representation of the attribute is merely a syntactic sugar and is converted to an integer attribute by the parser.
The cmpi operation is a generic comparison for integer-like types. Its two arguments can be integers, vectors or tensors thereof as long as their types match. The operation produces an i1 for the former case, a vector or a tensor of i1 with the same shape as inputs in the other cases.
Its first argument is an attribute that defines which type of comparison is performed. The following comparisons are supported:
equal (mnemonic: "eq"; integer value: 0)
not equal (mnemonic: "ne"; integer value: 1)
signed less than (mnemonic: "slt"; integer value: 2)
signed less than or equal (mnemonic: "sle"; integer value: 3)
signed greater than (mnemonic: "sgt"; integer value: 4)
signed greater than or equal (mnemonic: "sge"; integer value: 5)
unsigned less than (mnemonic: "ult"; integer value: 6)
unsigned less than or equal (mnemonic: "ule"; integer value: 7)
unsigned greater than (mnemonic: "ugt"; integer value: 8)
unsigned greater than or equal (mnemonic: "uge"; integer value: 9)
The result is 1 if the comparison is true and 0 otherwise. For vector or tensor operands, the comparison is performed elementwise and the element of the result indicates whether the comparison is true for the operand elements with the same indices as those of the result.
Note: while the custom assembly form uses strings, the actual underlying attribute has integer type (or rather enum class in C++ code) as seen from the generic assembly form. String literals are used to improve readability of the IR by humans.
This operation only applies to integer-like operands, but not floats. The main reason being that comparison operations have diverging sets of attributes: integers require sign specification while floats require various floating point-related particularities, e.g., -ffast-math behavior, IEEE754 compliance, etc (rationale). The type of comparison is specified as attribute to avoid introducing ten similar operations, taking into account that they are often implemented using the same operation downstream (rationale). The separation between signed and unsigned order comparisons is necessary because of integers being signless. The comparison operation must know how to interpret values with the foremost bit being set: negatives in two's complement or large positives (rationale).
Example
mlir
// Custom form of scalar "signed less than" comparison.
+%x = arith.cmpi slt, %lhs, %rhs : i32
+
+// Generic form of the same operation.
+%x = "arith.cmpi"(%lhs, %rhs) {predicate = 2 : i64} : (i32, i32) -> i1
+
+// Custom form of vector equality comparison.
+%x = arith.cmpi eq, %lhs, %rhs : vector<4xi64>
+
+// Generic form of the same operation.
+%x = "arith.cmpi"(%lhs, %rhs) {predicate = 0 : i64}
+ : (vector<4xi64>, vector<4xi64>) -> vector<4xi1>
The constant operation produces an SSA value equal to some integer or floating-point constant specified by an attribute. This is the way MLIR forms simple integer and floating point constants.
Signed integer division. Rounds towards zero. Treats the leading bit as sign, i.e. 6 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
Unsigned integer division. Rounds towards zero. Treats the leading bit as the most significant, i.e. for i16 given two's complement representation, 6 / -2 = 6 / (2^16 - 2) = 0.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
`,6))]),a("details",L,[a("summary",null,[e[39]||(e[39]=a("a",{id:"Reactant.MLIR.Dialects.arith.extf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.extf-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.extf")],-1)),e[40]||(e[40]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[41]||(e[41]=a("p",null,[a("code",null,"extf")],-1)),e[42]||(e[42]=a("p",null,"Cast a floating-point value to a larger floating-point-typed value. The destination type must to be strictly wider than the source type. When operating on vectors, casts elementwise.",-1)),e[43]||(e[43]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L632-L638",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",j,[a("summary",null,[e[44]||(e[44]=a("a",{id:"Reactant.MLIR.Dialects.arith.extsi-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.extsi-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.extsi")],-1)),e[45]||(e[45]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[46]||(e[46]=i(`
extsi
The integer sign extension operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be larger than the input bit-width (N > M). The top-most (N - M) bits of the output are filled with copies of the most-significant bit of the input.
Example
mlir
%1 = arith.constant 5 : i3 // %1 is 0b101
+%2 = arith.extsi %1 : i3 to i6 // %2 is 0b111101
+%3 = arith.constant 2 : i3 // %3 is 0b010
+%4 = arith.extsi %3 : i3 to i6 // %4 is 0b000010
+
+%5 = arith.extsi %0 : vector<2 x i32> to vector<2 x i64>
The integer zero extension operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be larger than the input bit-width (N > M). The top-most (N - M) bits of the output are filled with zeros.
Example
mlir
%1 = arith.constant 5 : i3 // %1 is 0b101
+ %2 = arith.extui %1 : i3 to i6 // %2 is 0b000101
+ %3 = arith.constant 2 : i3 // %3 is 0b010
+ %4 = arith.extui %3 : i3 to i6 // %4 is 0b000010
+
+ %5 = arith.extui %0 : vector<2 x i32> to vector<2 x i64>
Signed integer division. Rounds towards negative infinity, i.e. 5 / -2 = -3.
Divison by zero, or signed division overflow (minimum value divided by -1) 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.
`,6))]),a("details",T,[a("summary",null,[e[53]||(e[53]=a("a",{id:"Reactant.MLIR.Dialects.arith.fptosi-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.fptosi-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.fptosi")],-1)),e[54]||(e[54]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[55]||(e[55]=a("p",null,[a("code",null,"fptosi")],-1)),e[56]||(e[56]=a("p",null,"Cast from a value interpreted as floating-point to the nearest (rounding towards zero) signed integer value. When operating on vectors, casts elementwise.",-1)),e[57]||(e[57]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L736-L742",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",w,[a("summary",null,[e[58]||(e[58]=a("a",{id:"Reactant.MLIR.Dialects.arith.fptoui-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.fptoui-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.fptoui")],-1)),e[59]||(e[59]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[60]||(e[60]=a("p",null,[a("code",null,"fptoui")],-1)),e[61]||(e[61]=a("p",null,"Cast from a value interpreted as floating-point to the nearest (rounding towards zero) unsigned integer value. When operating on vectors, casts elementwise.",-1)),e[62]||(e[62]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L762-L768",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",k,[a("summary",null,[e[63]||(e[63]=a("a",{id:"Reactant.MLIR.Dialects.arith.index_cast-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.index_cast-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.index_cast")],-1)),e[64]||(e[64]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[65]||(e[65]=a("p",null,[a("code",null,"index_cast")],-1)),e[66]||(e[66]=a("p",null,"Casts between scalar or vector integers and corresponding 'index' scalar or vectors. Index is an integer of platform-specific bit width. If casting to a wider integer, the value is sign-extended. If casting to a narrower integer, the value is truncated.",-1)),e[67]||(e[67]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L828-L835",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",V,[a("summary",null,[e[68]||(e[68]=a("a",{id:"Reactant.MLIR.Dialects.arith.index_castui-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.index_castui-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.index_castui")],-1)),e[69]||(e[69]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[70]||(e[70]=a("p",null,[a("code",null,"index_castui")],-1)),e[71]||(e[71]=a("p",null,"Casts between scalar or vector integers and corresponding 'index' scalar or vectors. Index is an integer of platform-specific bit width. If casting to a wider integer, the value is zero-extended. If casting to a narrower integer, the value is truncated.",-1)),e[72]||(e[72]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L855-L862",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",A,[a("summary",null,[e[73]||(e[73]=a("a",{id:"Reactant.MLIR.Dialects.arith.maximumf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.maximumf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.maximumf")],-1)),e[74]||(e[74]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[75]||(e[75]=i(`
maximumf
Returns the maximum of the two arguments, treating -0.0 as less than +0.0. If one of the arguments is NaN, then the result is also NaN.
Returns the maximum of the two arguments. If the arguments are -0.0 and +0.0, then the result is either of them. If one of the arguments is NaN, then the result is the other argument.
Returns the minimum of the two arguments. If the arguments are -0.0 and +0.0, then the result is either of them. If one of the arguments is NaN, then the result is the other argument.
The mulf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit multiplication on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the multiplication modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The muli operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
Performs (2*N)-bit multiplication on sign-extended operands. Returns two N-bit results: the low and the high halves of the product. The low half has the same value as the result of regular multiplication arith.muli with the same operands.
Performs (2*N)-bit multiplication on zero-extended operands. Returns two N-bit results: the low and the high halves of the product. The low half has the same value as the result of regular multiplication arith.muli with the same operands.
The negf operation computes the negation of a given value. It takes one operand and returns one result of the same type. This type may be a float scalar type, a vector whose element type is float, or a tensor of floats. It has no standard attributes.
The ori operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
`,5))]),a("details",P,[a("summary",null,[e[103]||(e[103]=a("a",{id:"Reactant.MLIR.Dialects.arith.remf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.remf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.remf")],-1)),e[104]||(e[104]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[105]||(e[105]=a("p",null,[a("code",null,"remf")],-1)),e[106]||(e[106]=a("p",null,"Returns the floating point division remainder. The remainder has the same sign as the dividend (lhs operand).",-1)),e[107]||(e[107]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1431-L1436",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",U,[a("summary",null,[e[108]||(e[108]=a("a",{id:"Reactant.MLIR.Dialects.arith.remsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.remsi-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.remsi")],-1)),e[109]||(e[109]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[110]||(e[110]=i(`
remsi
Signed integer division remainder. Treats the leading bit as sign, i.e. 6 % -2 = 0.
Division by zero is undefined behavior. When applied to vector and tensor values, the behavior is undefined if any elements are divided by zero.
The arith.select operation chooses one value based on a binary condition supplied as its first operand.
If the value of the first operand (the condition) is 1, then the second operand is returned, and the third operand is ignored, even if it was poison.
If the value of the first operand (the condition) is 0, then the third operand is returned, and the second operand is ignored, even if it was poison.
If the value of the first operand (the condition) is poison, then the operation returns poison.
The operation applies to vectors and tensors elementwise given the shape of all operands is identical. The choice is made for each element individually based on the value at the same position as the element in the condition operand. If an i1 is provided as the condition, the entire vector or tensor is chosen.
Example
mlir
// Custom form of scalar selection.
+%x = arith.select %cond, %true, %false : i32
+
+// Generic form of the same operation.
+%x = "arith.select"(%cond, %true, %false) : (i1, i32, i32) -> i32
+
+// Element-wise vector selection.
+%vx = arith.select %vcond, %vtrue, %vfalse : vector<42xi1>, vector<42xf32>
+
+// Full vector selection.
+%vx = arith.select %cond, %vtrue, %vfalse : vector<42xf32>
The shli operation shifts the integer value of the first operand to the left by the integer value of the second operand. The second operand is interpreted as unsigned. The low order bits are 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.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
The shrsi operation shifts an integer value of the first operand to the right by the value of the second operand. The first operand is interpreted as signed, and the second operand is interpreter as unsigned. The high order bits in the output are filled with copies of the most-significant bit of the shifted value (which means that the sign of the value is preserved). If the value of the second operand is greater or equal than bitwidth of the first operand, then the operation returns poison.
The shrui operation shifts an integer value of the first operand to the right by the value of the second operand. The first operand is interpreted as unsigned, and the second operand is interpreted as unsigned. The high order bits are always 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.
`,5))]),a("details",K,[a("summary",null,[e[126]||(e[126]=a("a",{id:"Reactant.MLIR.Dialects.arith.sitofp-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.sitofp-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.sitofp")],-1)),e[127]||(e[127]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[128]||(e[128]=a("p",null,[a("code",null,"sitofp")],-1)),e[129]||(e[129]=a("p",null,"Cast from a value interpreted as a signed integer to the corresponding floating-point value. If the value cannot be exactly represented, it is rounded using the default rounding mode. When operating on vectors, casts elementwise.",-1)),e[130]||(e[130]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1554-L1561",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",Q,[a("summary",null,[e[131]||(e[131]=a("a",{id:"Reactant.MLIR.Dialects.arith.subf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.subf-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.subf")],-1)),e[132]||(e[132]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[133]||(e[133]=i(`
subf
The subf operation takes two operands and returns one result, each of these is required to be the same type. This type may be a floating point scalar type, a vector whose element type is a floating point type, or a floating point tensor.
Performs N-bit subtraction on the operands. The operands are interpreted as unsigned bitvectors. The result is represented by a bitvector containing the mathematical value of the subtraction modulo 2^n, where n is the bitwidth. Because arith integers use a two's complement representation, this operation is applicable on both signed and unsigned integer operands.
The subi operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers.
This op supports nuw/nsw overflow flags which stands stand for "No Unsigned Wrap" and "No Signed Wrap", respectively. If the nuw and/or nsw flags are present, and an unsigned/signed overflow occurs (respectively), the result is poison.
`,7))]),a("details",Y,[a("summary",null,[e[137]||(e[137]=a("a",{id:"Reactant.MLIR.Dialects.arith.truncf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.truncf-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.truncf")],-1)),e[138]||(e[138]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[139]||(e[139]=a("p",null,[a("code",null,"truncf")],-1)),e[140]||(e[140]=a("p",null,"Truncate a floating-point value to a smaller floating-point-typed value. The destination type must be strictly narrower than the source type. If the value cannot be exactly represented, it is rounded using the provided rounding mode or the default one if no rounding mode is provided. When operating on vectors, casts elementwise.",-1)),e[141]||(e[141]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1827-L1835",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",Z,[a("summary",null,[e[142]||(e[142]=a("a",{id:"Reactant.MLIR.Dialects.arith.trunci-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.trunci-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.trunci")],-1)),e[143]||(e[143]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[144]||(e[144]=i(`
trunci
The integer truncation operation takes an integer input of width M and an integer destination type of width N. The destination bit-width must be smaller than the input bit-width (N < M). The top-most (N - M) bits of the input are discarded.
Example
mlir
%1 = arith.constant 21 : i5 // %1 is 0b10101
+ %2 = arith.trunci %1 : i5 to i4 // %2 is 0b0101
+ %3 = arith.trunci %1 : i5 to i3 // %3 is 0b101
+
+ %5 = arith.trunci %0 : vector<2 x i32> to vector<2 x i16>
`,5))]),a("details",_,[a("summary",null,[e[145]||(e[145]=a("a",{id:"Reactant.MLIR.Dialects.arith.uitofp-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.uitofp-Tuple{Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.uitofp")],-1)),e[146]||(e[146]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[147]||(e[147]=a("p",null,[a("code",null,"uitofp")],-1)),e[148]||(e[148]=a("p",null,"Cast from a value interpreted as unsigned integer to the corresponding floating-point value. If the value cannot be exactly represented, it is rounded using the default rounding mode. When operating on vectors, casts elementwise.",-1)),e[149]||(e[149]=a("p",null,[a("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Arith.jl#L1897-L1904",target:"_blank",rel:"noreferrer"},"source")],-1))]),a("details",ee,[a("summary",null,[e[150]||(e[150]=a("a",{id:"Reactant.MLIR.Dialects.arith.xori-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.arith.xori-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[a("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.arith.xori")],-1)),e[151]||(e[151]=t()),n(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[152]||(e[152]=i(`
xori
The xori operation takes two operands and returns one result, each of these is required to be the same type. This type may be an integer scalar type, a vector whose element type is integer, or a tensor of integers. It has no standard attributes.
`,5))])])}const re=l(c,[["render",ae]]);export{pe as __pageData,re as default};
diff --git a/previews/PR363/assets/api_builtin.md.XATQzjrL.js b/previews/PR363/assets/api_builtin.md.XATQzjrL.js
new file mode 100644
index 00000000..ee5779bc
--- /dev/null
+++ b/previews/PR363/assets/api_builtin.md.XATQzjrL.js
@@ -0,0 +1,18 @@
+import{_ as i,c as l,j as n,a,G as t,a2 as o,B as p,o as r}from"./chunks/framework.2yyKLD8d.js";const _=JSON.parse('{"title":"Builtin Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/builtin.md","filePath":"api/builtin.md","lastUpdated":null}'),c={name:"api/builtin.md"},d={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"};function m(f,e,b,h,y,g){const s=p("Badge");return r(),l("div",null,[e[6]||(e[6]=n("h1",{id:"Builtin-Dialect",tabindex:"-1"},[a("Builtin Dialect "),n("a",{class:"header-anchor",href:"#Builtin-Dialect","aria-label":'Permalink to "Builtin Dialect {#Builtin-Dialect}"'},"")],-1)),e[7]||(e[7]=n("p",null,[a("Refer to the "),n("a",{href:"https://mlir.llvm.org/docs/Dialects/Builtin/",target:"_blank",rel:"noreferrer"},"official documentation"),a(" for more details.")],-1)),n("details",d,[n("summary",null,[e[0]||(e[0]=n("a",{id:"Reactant.MLIR.Dialects.builtin.module_-Tuple{}",href:"#Reactant.MLIR.Dialects.builtin.module_-Tuple{}"},[n("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.builtin.module_")],-1)),e[1]||(e[1]=a()),t(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=o(`
module_
A module represents a top-level container operation. It contains a single graph region containing a single block which can contain any operations and does not have a terminator. Operations within this region cannot implicitly capture values defined outside the module, i.e. Modules are IsolatedFromAbove. Modules have an optional symbol name which can be used to refer to them in operations.
An unrealized_conversion_cast operation represents an unrealized conversion from one set of types to another, that is used to enable the inter-mixing of different type systems. This operation should not be attributed any special representational or execution semantics, and is generally only intended to be used to satisfy the temporary intermixing of type systems during the conversion of one type system to another.
This operation may produce results of arity 1-N, and accept as input operands of arity 0-N.
Example
mlir
// An unrealized 0-1 conversion. These types of conversions are useful in
+// cases where a type is removed from the type system, but not all uses have
+// been converted. For example, imagine we have a tuple type that is
+// expanded to its element types. If only some uses of an empty tuple type
+// instance are converted we still need an instance of the tuple type, but
+// have no inputs to the unrealized conversion.
+%result = unrealized_conversion_cast to !bar.tuple_type<>
+
+// An unrealized 1-1 conversion.
+%result1 = unrealized_conversion_cast %operand : !foo.type to !bar.lowered_type
+
+// An unrealized 1-N conversion.
+%results2:2 = unrealized_conversion_cast %tuple_operand : !foo.tuple_type<!foo.type, !foo.type> to !foo.type, !foo.type
+
+// An unrealized N-1 conversion.
+%result3 = unrealized_conversion_cast %operand, %operand : !foo.type, !foo.type to !bar.tuple_type<!foo.type, !foo.type>
`,6))])])}const R=i(c,[["render",m]]);export{_ as __pageData,R as default};
diff --git a/previews/PR363/assets/api_builtin.md.XATQzjrL.lean.js b/previews/PR363/assets/api_builtin.md.XATQzjrL.lean.js
new file mode 100644
index 00000000..ee5779bc
--- /dev/null
+++ b/previews/PR363/assets/api_builtin.md.XATQzjrL.lean.js
@@ -0,0 +1,18 @@
+import{_ as i,c as l,j as n,a,G as t,a2 as o,B as p,o as r}from"./chunks/framework.2yyKLD8d.js";const _=JSON.parse('{"title":"Builtin Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/builtin.md","filePath":"api/builtin.md","lastUpdated":null}'),c={name:"api/builtin.md"},d={class:"jldocstring custom-block"},u={class:"jldocstring custom-block"};function m(f,e,b,h,y,g){const s=p("Badge");return r(),l("div",null,[e[6]||(e[6]=n("h1",{id:"Builtin-Dialect",tabindex:"-1"},[a("Builtin Dialect "),n("a",{class:"header-anchor",href:"#Builtin-Dialect","aria-label":'Permalink to "Builtin Dialect {#Builtin-Dialect}"'},"")],-1)),e[7]||(e[7]=n("p",null,[a("Refer to the "),n("a",{href:"https://mlir.llvm.org/docs/Dialects/Builtin/",target:"_blank",rel:"noreferrer"},"official documentation"),a(" for more details.")],-1)),n("details",d,[n("summary",null,[e[0]||(e[0]=n("a",{id:"Reactant.MLIR.Dialects.builtin.module_-Tuple{}",href:"#Reactant.MLIR.Dialects.builtin.module_-Tuple{}"},[n("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.builtin.module_")],-1)),e[1]||(e[1]=a()),t(s,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=o(`
module_
A module represents a top-level container operation. It contains a single graph region containing a single block which can contain any operations and does not have a terminator. Operations within this region cannot implicitly capture values defined outside the module, i.e. Modules are IsolatedFromAbove. Modules have an optional symbol name which can be used to refer to them in operations.
An unrealized_conversion_cast operation represents an unrealized conversion from one set of types to another, that is used to enable the inter-mixing of different type systems. This operation should not be attributed any special representational or execution semantics, and is generally only intended to be used to satisfy the temporary intermixing of type systems during the conversion of one type system to another.
This operation may produce results of arity 1-N, and accept as input operands of arity 0-N.
Example
mlir
// An unrealized 0-1 conversion. These types of conversions are useful in
+// cases where a type is removed from the type system, but not all uses have
+// been converted. For example, imagine we have a tuple type that is
+// expanded to its element types. If only some uses of an empty tuple type
+// instance are converted we still need an instance of the tuple type, but
+// have no inputs to the unrealized conversion.
+%result = unrealized_conversion_cast to !bar.tuple_type<>
+
+// An unrealized 1-1 conversion.
+%result1 = unrealized_conversion_cast %operand : !foo.type to !bar.lowered_type
+
+// An unrealized 1-N conversion.
+%results2:2 = unrealized_conversion_cast %tuple_operand : !foo.tuple_type<!foo.type, !foo.type> to !foo.type, !foo.type
+
+// An unrealized N-1 conversion.
+%result3 = unrealized_conversion_cast %operand, %operand : !foo.type, !foo.type to !bar.tuple_type<!foo.type, !foo.type>
`,6))])])}const R=i(c,[["render",m]]);export{_ as __pageData,R as default};
diff --git a/previews/PR363/assets/api_chlo.md.CTmmK6a-.js b/previews/PR363/assets/api_chlo.md.CTmmK6a-.js
new file mode 100644
index 00000000..890c9f99
--- /dev/null
+++ b/previews/PR363/assets/api_chlo.md.CTmmK6a-.js
@@ -0,0 +1 @@
+import{_ as o,c as r,j as l,a as e,G as n,a2 as s,B as i,o as u}from"./chunks/framework.2yyKLD8d.js";const Rt=JSON.parse('{"title":"CHLO Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/chlo.md","filePath":"api/chlo.md","lastUpdated":null}'),d={name:"api/chlo.md"},R={class:"jldocstring custom-block"},p={class:"jldocstring custom-block"},c={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"},I={class:"jldocstring custom-block"},L={class:"jldocstring custom-block"},M={class:"jldocstring custom-block"},j={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},D={class:"jldocstring custom-block"},w={class:"jldocstring custom-block"},x={class:"jldocstring custom-block"},T={class:"jldocstring custom-block"},h={class:"jldocstring custom-block"},V={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"},O={class:"jldocstring custom-block"},z={class:"jldocstring custom-block"},C={class:"jldocstring custom-block"},A={class:"jldocstring custom-block"},E={class:"jldocstring custom-block"},H={class:"jldocstring custom-block"},v={class:"jldocstring custom-block"},$={class:"jldocstring custom-block"},S={class:"jldocstring custom-block"},q={class:"jldocstring custom-block"},N={class:"jldocstring custom-block"},B={class:"jldocstring custom-block"},P={class:"jldocstring custom-block"},G={class:"jldocstring custom-block"},F={class:"jldocstring custom-block"},U={class:"jldocstring custom-block"},Z={class:"jldocstring custom-block"},J={class:"jldocstring custom-block"},K={class:"jldocstring custom-block"},Q={class:"jldocstring custom-block"},W={class:"jldocstring custom-block"},X={class:"jldocstring custom-block"},Y={class:"jldocstring custom-block"},_={class:"jldocstring custom-block"},tt={class:"jldocstring custom-block"},lt={class:"jldocstring custom-block"},et={class:"jldocstring custom-block"},at={class:"jldocstring custom-block"};function nt(st,t,ot,rt,it,ut){const a=i("Badge");return u(),r("div",null,[t[277]||(t[277]=l("h1",{id:"CHLO-Dialect",tabindex:"-1"},[e("CHLO Dialect "),l("a",{class:"header-anchor",href:"#CHLO-Dialect","aria-label":'Permalink to "CHLO Dialect {#CHLO-Dialect}"'},"")],-1)),t[278]||(t[278]=l("p",null,[e("Refer to the "),l("a",{href:"https://github.com/openxla/xla/tree/main/xla/mlir_hlo#hlo-client-dialect-chlo",target:"_blank",rel:"noreferrer"},"official documentation"),e(" for more details.")],-1)),l("details",R,[l("summary",null,[t[0]||(t[0]=l("a",{id:"Reactant.MLIR.Dialects.chlo._asin_acos_kernel-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo._asin_acos_kernel-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo._asin_acos_kernel")],-1)),t[1]||(t[1]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[2]||(t[2]=l("p",null,[l("code",null,"_asin_acos_kernel")],-1)),t[3]||(t[3]=l("p",null,[e("Returns "),l("code",null,"AsinAcosKernel(operand)"),e(" element-wise.")],-1)),t[4]||(t[4]=l("p",null,"If w = _asin_acos_kernel(z) w' = _asin_acos_kernel(I * z) then asin(z) = complex(atan2(z.real, w.real), sign(z.imag) * w.imag) acos(z) = complex(atan2(w.real, z.real), -sign(z.imag) * w.imag) asinh(z) = complex(sign(z.real) * w'.imag, atan2(z.imag, w'.real)) acosh(z) = complex(w.imag, sign(z.imag) * atan2(w.real, z.real))",-1)),t[5]||(t[5]=l("p",null,"This op is used as an intermediate value in decompositions and should never be constructed directly by frameworks or consumed by backends.",-1)),t[6]||(t[6]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L76-L93",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",p,[l("summary",null,[t[7]||(t[7]=l("a",{id:"Reactant.MLIR.Dialects.chlo.acos-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.acos-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.acos")],-1)),t[8]||(t[8]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[9]||(t[9]=l("p",null,[l("code",null,"acos")],-1)),t[10]||(t[10]=l("p",null,[e("Returns "),l("code",null,"Acos(operand)"),e(" element-wise.")],-1)),t[11]||(t[11]=l("p",null,"$",-1)),t[12]||(t[12]=l("p",null,"\\acos(x) = 2 * \\atan(\\sqrt(1 - x^2) / (1 + x)) if x != -1 = pi if x == -1 $",-1)),t[13]||(t[13]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L16-L25",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",c,[l("summary",null,[t[14]||(t[14]=l("a",{id:"Reactant.MLIR.Dialects.chlo.acosh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.acosh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.acosh")],-1)),t[15]||(t[15]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[16]||(t[16]=l("p",null,[l("code",null,"acosh")],-1)),t[17]||(t[17]=l("p",null,[e("Returns "),l("code",null,"Acosh(operand)"),e(" element-wise.")],-1)),t[18]||(t[18]=l("p",null,"$",-1)),t[19]||(t[19]=l("p",null,"\\acosh(x) = log(x + sqrt(x^2 - 1)) if x >= -1 \\acosh(x) = nan if x < -1 $",-1)),t[20]||(t[20]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L46-L55",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",m,[l("summary",null,[t[21]||(t[21]=l("a",{id:"Reactant.MLIR.Dialects.chlo.asin-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.asin-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.asin")],-1)),t[22]||(t[22]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[23]||(t[23]=l("p",null,[l("code",null,"asin")],-1)),t[24]||(t[24]=l("p",null,[e("Returns "),l("code",null,"Asin(operand)"),e(" element-wise.")],-1)),t[25]||(t[25]=l("p",null,"$",-1)),t[26]||(t[26]=l("p",null,"\\asin(x) = 2 * atan(x / (1 + sqrt(1 - x^2))) $",-1)),t[27]||(t[27]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L116-L124",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",f,[l("summary",null,[t[28]||(t[28]=l("a",{id:"Reactant.MLIR.Dialects.chlo.asinh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.asinh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.asinh")],-1)),t[29]||(t[29]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[30]||(t[30]=l("p",null,[l("code",null,"asinh")],-1)),t[31]||(t[31]=l("p",null,[e("Returns "),l("code",null,"Asinh(operand)"),e(" element-wise.")],-1)),t[32]||(t[32]=l("p",null,"$",-1)),t[33]||(t[33]=l("p",null,"\\asinh(x) = log(x + sqrt(x^2 + 1)) $",-1)),t[34]||(t[34]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L145-L153",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",b,[l("summary",null,[t[35]||(t[35]=l("a",{id:"Reactant.MLIR.Dialects.chlo.atan-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.atan-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.atan")],-1)),t[36]||(t[36]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[37]||(t[37]=l("p",null,[l("code",null,"atan")],-1)),t[38]||(t[38]=l("p",null,[e("Returns "),l("code",null,"Atan(operand)"),e(" element-wise.")],-1)),t[39]||(t[39]=l("p",null,"$",-1)),t[40]||(t[40]=l("p",null,"\\atan(x) = \\atan2(x, 1) $",-1)),t[41]||(t[41]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L174-L182",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",I,[l("summary",null,[t[42]||(t[42]=l("a",{id:"Reactant.MLIR.Dialects.chlo.atanh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.atanh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.atanh")],-1)),t[43]||(t[43]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[44]||(t[44]=l("p",null,[l("code",null,"atanh")],-1)),t[45]||(t[45]=l("p",null,[e("Returns "),l("code",null,"Atanh(operand)"),e(" element-wise.")],-1)),t[46]||(t[46]=l("p",null,"$",-1)),t[47]||(t[47]=l("p",null,"\\atanh(x) = 0.5 * log((1 + x) / (1 - x)) if abs(x) <= 1 = nan otherwise $",-1)),t[48]||(t[48]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L203-L212",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",L,[l("summary",null,[t[49]||(t[49]=l("a",{id:"Reactant.MLIR.Dialects.chlo.bessel_i1e-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.bessel_i1e-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.bessel_i1e")],-1)),t[50]||(t[50]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[51]||(t[51]=l("p",null,[l("code",null,"bessel_i1e")],-1)),t[52]||(t[52]=l("p",null,[e("Returns "),l("code",null,"bessel_i1e(operand)"),e(" element-wise.")],-1)),t[53]||(t[53]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L233-L237",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",M,[l("summary",null,[t[54]||(t[54]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_add-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_add-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_add")],-1)),t[55]||(t[55]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[56]||(t[56]=l("p",null,[l("code",null,"broadcast_add")],-1)),t[57]||(t[57]=l("p",null,[e("Returns "),l("code",null,"lhs + rhs"),e(" element-wise.")],-1)),t[58]||(t[58]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[59]||(t[59]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L260-L267",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",j,[l("summary",null,[t[60]||(t[60]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_and-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_and-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_and")],-1)),t[61]||(t[61]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[62]||(t[62]=l("p",null,[l("code",null,"broadcast_and")],-1)),t[63]||(t[63]=l("p",null,[e("Returns "),l("code",null,"logical_and(lhs, rhs)"),e(" element-wise.")],-1)),t[64]||(t[64]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[65]||(t[65]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L296-L303",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",g,[l("summary",null,[t[66]||(t[66]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_atan2-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_atan2-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_atan2")],-1)),t[67]||(t[67]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[68]||(t[68]=l("p",null,[l("code",null,"broadcast_atan2")],-1)),t[69]||(t[69]=l("p",null,[e("Returns "),l("code",null,"atan2(lhs/rhs)"),e(" element-wise.")],-1)),t[70]||(t[70]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[71]||(t[71]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L332-L339",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",y,[l("summary",null,[t[72]||(t[72]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_compare-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_compare-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_compare")],-1)),t[73]||(t[73]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[74]||(t[74]=s('
broadcast_compare
Compares lhs and rhs elementwise according to comparison_direction and compare_type. If unspecified, compare_type is FLOAT for float element types, SIGNED for signed element types and UNSIGNED for unsigned element types.
',4))]),l("details",D,[l("summary",null,[t[75]||(t[75]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_complex-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_complex-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_complex")],-1)),t[76]||(t[76]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[77]||(t[77]=l("p",null,[l("code",null,"broadcast_complex")],-1)),t[78]||(t[78]=l("p",null,"Performs element-wise conversion of a pair of real and imaginary values to a complex value.",-1)),t[79]||(t[79]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L413-L418",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",w,[l("summary",null,[t[80]||(t[80]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_divide-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_divide-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_divide")],-1)),t[81]||(t[81]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[82]||(t[82]=l("p",null,[l("code",null,"broadcast_divide")],-1)),t[83]||(t[83]=l("p",null,[e("Returns "),l("code",null,"lhs / rhs"),e(" element-wise.")],-1)),t[84]||(t[84]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[85]||(t[85]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L447-L454",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",x,[l("summary",null,[t[86]||(t[86]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_maximum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_maximum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_maximum")],-1)),t[87]||(t[87]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[88]||(t[88]=l("p",null,[l("code",null,"broadcast_maximum")],-1)),t[89]||(t[89]=l("p",null,[e("Returns "),l("code",null,"max(lhs, rhs)"),e(" element-wise.")],-1)),t[90]||(t[90]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[91]||(t[91]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L483-L490",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",T,[l("summary",null,[t[92]||(t[92]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_minimum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_minimum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_minimum")],-1)),t[93]||(t[93]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[94]||(t[94]=l("p",null,[l("code",null,"broadcast_minimum")],-1)),t[95]||(t[95]=l("p",null,[e("Returns "),l("code",null,"min(lhs, rhs)"),e(" element-wise.")],-1)),t[96]||(t[96]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[97]||(t[97]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L519-L526",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",h,[l("summary",null,[t[98]||(t[98]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_multiply-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_multiply-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_multiply")],-1)),t[99]||(t[99]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[100]||(t[100]=l("p",null,[l("code",null,"broadcast_multiply")],-1)),t[101]||(t[101]=l("p",null,[e("Returns "),l("code",null,"lhs * rhs"),e(" element-wise.")],-1)),t[102]||(t[102]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[103]||(t[103]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L555-L562",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",V,[l("summary",null,[t[104]||(t[104]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_next_after")],-1)),t[105]||(t[105]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[106]||(t[106]=l("p",null,[l("code",null,"broadcast_next_after")],-1)),t[107]||(t[107]=l("p",null,[e("Returns the next representable value of "),l("code",null,"lhs"),e(" in the direction of "),l("code",null,"rhs"),e(", element-wise. It can also return a subnormal number.")],-1)),t[108]||(t[108]=l("p",null,"Equivalent to the C++ std::nextafter function.",-1)),t[109]||(t[109]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L591-L598",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",k,[l("summary",null,[t[110]||(t[110]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_or-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_or-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_or")],-1)),t[111]||(t[111]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[112]||(t[112]=l("p",null,[l("code",null,"broadcast_or")],-1)),t[113]||(t[113]=l("p",null,[e("Returns "),l("code",null,"logical_or(lhs, rhs)"),e(" element-wise.")],-1)),t[114]||(t[114]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[115]||(t[115]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L627-L634",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",O,[l("summary",null,[t[116]||(t[116]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_polygamma")],-1)),t[117]||(t[117]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[118]||(t[118]=l("p",null,[l("code",null,"broadcast_polygamma")],-1)),t[119]||(t[119]=l("p",null,[e("Returns "),l("code",null,"Polygamma(operand, operand)"),e(" element-wise.")],-1)),t[120]||(t[120]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L663-L667",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",z,[l("summary",null,[t[121]||(t[121]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_power-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_power-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_power")],-1)),t[122]||(t[122]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[123]||(t[123]=l("p",null,[l("code",null,"broadcast_power")],-1)),t[124]||(t[124]=l("p",null,[e("Returns "),l("code",null,"lhs ^ rhs"),e(" element-wise.")],-1)),t[125]||(t[125]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[126]||(t[126]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L696-L703",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",C,[l("summary",null,[t[127]||(t[127]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_remainder-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_remainder-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_remainder")],-1)),t[128]||(t[128]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[129]||(t[129]=l("p",null,[l("code",null,"broadcast_remainder")],-1)),t[130]||(t[130]=l("p",null,[e("Returns "),l("code",null,"lhs % rhs"),e(" element-wise.")],-1)),t[131]||(t[131]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[132]||(t[132]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L732-L739",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",A,[l("summary",null,[t[133]||(t[133]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_select-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_select-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_select")],-1)),t[134]||(t[134]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[135]||(t[135]=l("p",null,[l("code",null,"broadcast_select")],-1)),t[136]||(t[136]=l("p",null,"Constructs an output array from elements of two input arrays, based on the values of a predicate array.",-1)),t[137]||(t[137]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#select",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#select")],-1)),t[138]||(t[138]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L768-L775",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",E,[l("summary",null,[t[139]||(t[139]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_shift_left-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_shift_left-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_shift_left")],-1)),t[140]||(t[140]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[141]||(t[141]=l("p",null,[l("code",null,"broadcast_shift_left")],-1)),t[142]||(t[142]=l("p",null,[e("Returns "),l("code",null,"lhs << rhs"),e(" element-wise.")],-1)),t[143]||(t[143]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[144]||(t[144]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L802-L809",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",H,[l("summary",null,[t[145]||(t[145]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_arithmetic-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_shift_right_arithmetic-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_arithmetic")],-1)),t[146]||(t[146]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[147]||(t[147]=l("p",null,[l("code",null,"broadcast_shift_right_arithmetic")],-1)),t[148]||(t[148]=l("p",null,[e("Returns "),l("code",null,"lhs >> rhs"),e(" element-wise.")],-1)),t[149]||(t[149]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[150]||(t[150]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L838-L845",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",v,[l("summary",null,[t[151]||(t[151]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_logical-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_shift_right_logical-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_logical")],-1)),t[152]||(t[152]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[153]||(t[153]=l("p",null,[l("code",null,"broadcast_shift_right_logical")],-1)),t[154]||(t[154]=l("p",null,[e("Returns "),l("code",null,"lhs >> rhs"),e(" element-wise.")],-1)),t[155]||(t[155]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[156]||(t[156]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L874-L881",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",$,[l("summary",null,[t[157]||(t[157]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_subtract-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_subtract-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_subtract")],-1)),t[158]||(t[158]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[159]||(t[159]=l("p",null,[l("code",null,"broadcast_subtract")],-1)),t[160]||(t[160]=l("p",null,[e("Returns "),l("code",null,"lhs - rhs"),e(" element-wise.")],-1)),t[161]||(t[161]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[162]||(t[162]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L910-L917",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",S,[l("summary",null,[t[163]||(t[163]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_xor-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_xor-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_xor")],-1)),t[164]||(t[164]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[165]||(t[165]=l("p",null,[l("code",null,"broadcast_xor")],-1)),t[166]||(t[166]=l("p",null,[e("Returns "),l("code",null,"logical_xor(lhs, rhs)"),e(" element-wise.")],-1)),t[167]||(t[167]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[168]||(t[168]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L946-L953",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",q,[l("summary",null,[t[169]||(t[169]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_zeta")],-1)),t[170]||(t[170]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[171]||(t[171]=l("p",null,[l("code",null,"broadcast_zeta")],-1)),t[172]||(t[172]=l("p",null,[e("Returns "),l("code",null,"Zeta(operand, operand)"),e(" element-wise.")],-1)),t[173]||(t[173]=l("p",null,"$",-1)),t[174]||(t[174]=l("p",null,"(\\zeta(x, q) = \\sum_{n=0}^{\\infty} (q + n)^{-x}) $",-1)),t[175]||(t[175]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L982-L990",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",N,[l("summary",null,[t[176]||(t[176]=l("a",{id:"Reactant.MLIR.Dialects.chlo.conj-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.conj-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.conj")],-1)),t[177]||(t[177]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[178]||(t[178]=l("p",null,[l("code",null,"conj")],-1)),t[179]||(t[179]=l("p",null,[e("Returns "),l("code",null,"Conj(operand)"),e(" element-wise.")],-1)),t[180]||(t[180]=l("p",null,"$",-1)),t[181]||(t[181]=l("p",null,"\\conj(x) = (\\real(x), \\neg(\\imag(x))) $",-1)),t[182]||(t[182]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1019-L1027",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",B,[l("summary",null,[t[183]||(t[183]=l("a",{id:"Reactant.MLIR.Dialects.chlo.constant-Tuple{}",href:"#Reactant.MLIR.Dialects.chlo.constant-Tuple{}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.constant")],-1)),t[184]||(t[184]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[185]||(t[185]=l("p",null,[l("code",null,"constant")],-1)),t[186]||(t[186]=l("p",null,"Represents a constant value.",-1)),t[187]||(t[187]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1075-L1079",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",P,[l("summary",null,[t[188]||(t[188]=l("a",{id:"Reactant.MLIR.Dialects.chlo.constant_like-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.constant_like-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.constant_like")],-1)),t[189]||(t[189]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[190]||(t[190]=l("p",null,[l("code",null,"constant_like")],-1)),t[191]||(t[191]=l("p",null,"Returns a splat constant of the same shape as the operand.",-1)),t[192]||(t[192]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1048-L1052",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",G,[l("summary",null,[t[193]||(t[193]=l("a",{id:"Reactant.MLIR.Dialects.chlo.cosh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.cosh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.cosh")],-1)),t[194]||(t[194]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[195]||(t[195]=l("p",null,[l("code",null,"cosh")],-1)),t[196]||(t[196]=l("p",null,[e("Returns "),l("code",null,"Cosh(operand)"),e(" element-wise.")],-1)),t[197]||(t[197]=l("p",null,"$",-1)),t[198]||(t[198]=l("p",null,"\\cosh(x) = (e^x + e^-x) / 2 $",-1)),t[199]||(t[199]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1100-L1108",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",F,[l("summary",null,[t[200]||(t[200]=l("a",{id:"Reactant.MLIR.Dialects.chlo.digamma-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.digamma-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.digamma")],-1)),t[201]||(t[201]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[202]||(t[202]=l("p",null,[l("code",null,"digamma")],-1)),t[203]||(t[203]=l("p",null,[e("Returns "),l("code",null,"Digamma(operand)"),e(" element-wise.")],-1)),t[204]||(t[204]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1129-L1133",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",U,[l("summary",null,[t[205]||(t[205]=l("a",{id:"Reactant.MLIR.Dialects.chlo.erf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.erf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.erf")],-1)),t[206]||(t[206]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[207]||(t[207]=l("p",null,[l("code",null,"erf")],-1)),t[208]||(t[208]=l("p",null,[e("Computes the Gauss error function of "),l("code",null,"x"),e(" element-wise.")],-1)),t[209]||(t[209]=l("p",null,"erf(x) = erf_impl(x) if |x| < 1 = 1 - erfc_impl(x) otherwise",-1)),t[210]||(t[210]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1183-L1190",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",Z,[l("summary",null,[t[211]||(t[211]=l("a",{id:"Reactant.MLIR.Dialects.chlo.erf_inv-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.erf_inv-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.erf_inv")],-1)),t[212]||(t[212]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[213]||(t[213]=l("p",null,[l("code",null,"erf_inv")],-1)),t[214]||(t[214]=l("p",null,[e("Returns "),l("code",null,"ErfInv(operand)"),e(" element-wise.")],-1)),t[215]||(t[215]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1156-L1160",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",J,[l("summary",null,[t[216]||(t[216]=l("a",{id:"Reactant.MLIR.Dialects.chlo.erfc-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.erfc-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.erfc")],-1)),t[217]||(t[217]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[218]||(t[218]=l("p",null,[l("code",null,"erfc")],-1)),t[219]||(t[219]=l("p",null,"Computes an approximation of the error function complement (1 - erf(x)).",-1)),t[220]||(t[220]=l("p",null,"erfc(x) = erfc_impl(x) if |x| > 1 = 1 - erf_impl(x) otherwise",-1)),t[221]||(t[221]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1211-L1218",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",K,[l("summary",null,[t[222]||(t[222]=l("a",{id:"Reactant.MLIR.Dialects.chlo.is_inf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.is_inf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.is_inf")],-1)),t[223]||(t[223]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[224]||(t[224]=l("p",null,[l("code",null,"is_inf")],-1)),t[225]||(t[225]=l("p",null,"Returns if a value is +/-inf element-wise.",-1)),t[226]||(t[226]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1239-L1243",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",Q,[l("summary",null,[t[227]||(t[227]=l("a",{id:"Reactant.MLIR.Dialects.chlo.is_neg_inf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.is_neg_inf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.is_neg_inf")],-1)),t[228]||(t[228]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[229]||(t[229]=l("p",null,[l("code",null,"is_neg_inf")],-1)),t[230]||(t[230]=l("p",null,"Returns if a value is -inf element-wise.",-1)),t[231]||(t[231]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1264-L1268",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",W,[l("summary",null,[t[232]||(t[232]=l("a",{id:"Reactant.MLIR.Dialects.chlo.is_pos_inf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.is_pos_inf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.is_pos_inf")],-1)),t[233]||(t[233]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[234]||(t[234]=l("p",null,[l("code",null,"is_pos_inf")],-1)),t[235]||(t[235]=l("p",null,"Returns if a value is +inf element-wise.",-1)),t[236]||(t[236]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1291-L1295",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",X,[l("summary",null,[t[237]||(t[237]=l("a",{id:"Reactant.MLIR.Dialects.chlo.lgamma-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.lgamma-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.lgamma")],-1)),t[238]||(t[238]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[239]||(t[239]=l("p",null,[l("code",null,"lgamma")],-1)),t[240]||(t[240]=l("p",null,[e("Returns "),l("code",null,"Lgamma(operand)"),e(" element-wise.")],-1)),t[241]||(t[241]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1318-L1322",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",Y,[l("summary",null,[t[242]||(t[242]=l("a",{id:"Reactant.MLIR.Dialects.chlo.next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.next_after")],-1)),t[243]||(t[243]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[244]||(t[244]=l("p",null,[l("code",null,"next_after")],-1)),t[245]||(t[245]=l("p",null,[e("Returns the next representable value of "),l("code",null,"x"),e(" in the direction of "),l("code",null,"y"),e(", element-wise. It can also return a subnormal number.")],-1)),t[246]||(t[246]=l("p",null,"Equivalent to the C++ std::nextafter function.",-1)),t[247]||(t[247]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1343-L1350",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",_,[l("summary",null,[t[248]||(t[248]=l("a",{id:"Reactant.MLIR.Dialects.chlo.polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.polygamma")],-1)),t[249]||(t[249]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[250]||(t[250]=l("p",null,[l("code",null,"polygamma")],-1)),t[251]||(t[251]=l("p",null,[e("Returns "),l("code",null,"Polygamma(operand, operand)"),e(" element-wise.")],-1)),t[252]||(t[252]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1373-L1377",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",tt,[l("summary",null,[t[253]||(t[253]=l("a",{id:"Reactant.MLIR.Dialects.chlo.sinh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.sinh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.sinh")],-1)),t[254]||(t[254]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[255]||(t[255]=l("p",null,[l("code",null,"sinh")],-1)),t[256]||(t[256]=l("p",null,[e("Returns "),l("code",null,"Sinh(operand)"),e(" element-wise.")],-1)),t[257]||(t[257]=l("p",null,"$",-1)),t[258]||(t[258]=l("p",null,"\\sinh(x) = (e^x - e^-x) / 2 if |x| < 1 = e^(x + log(1/2)) - e^(-x + log(1/2)) otherwise. $",-1)),t[259]||(t[259]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1400-L1409",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",lt,[l("summary",null,[t[260]||(t[260]=l("a",{id:"Reactant.MLIR.Dialects.chlo.tan-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.tan-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.tan")],-1)),t[261]||(t[261]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[262]||(t[262]=l("p",null,[l("code",null,"tan")],-1)),t[263]||(t[263]=l("p",null,[e("Returns "),l("code",null,"Tan(operand)"),e(" element-wise.")],-1)),t[264]||(t[264]=l("p",null,"$",-1)),t[265]||(t[265]=l("p",null,"\\tan(x) = \\sin(x) / \\cos(x) $",-1)),t[266]||(t[266]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1430-L1438",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",et,[l("summary",null,[t[267]||(t[267]=l("a",{id:"Reactant.MLIR.Dialects.chlo.top_k-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.top_k-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.top_k")],-1)),t[268]||(t[268]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[269]||(t[269]=s('
top_k
If the input is a vector (rank-1), finds the k largest entries in the vector and outputs their values and indices as vectors. Thus values[j] is the j-th largest entry in input, and its index is indices[j].
For matrices (resp. higher rank input), computes the top k entries in each row (resp. vector along the last dimension). Thus,
',6))]),l("details",at,[l("summary",null,[t[270]||(t[270]=l("a",{id:"Reactant.MLIR.Dialects.chlo.zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.zeta")],-1)),t[271]||(t[271]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[272]||(t[272]=l("p",null,[l("code",null,"zeta")],-1)),t[273]||(t[273]=l("p",null,[e("Returns "),l("code",null,"Zeta(operand, operand)"),e(" element-wise.")],-1)),t[274]||(t[274]=l("p",null,"$",-1)),t[275]||(t[275]=l("p",null,"(\\zeta(x, q) = \\sum_{n=0}^{\\infty} (q + n)^{-x}) $",-1)),t[276]||(t[276]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1500-L1508",target:"_blank",rel:"noreferrer"},"source")],-1))])])}const pt=o(d,[["render",nt]]);export{Rt as __pageData,pt as default};
diff --git a/previews/PR363/assets/api_chlo.md.CTmmK6a-.lean.js b/previews/PR363/assets/api_chlo.md.CTmmK6a-.lean.js
new file mode 100644
index 00000000..890c9f99
--- /dev/null
+++ b/previews/PR363/assets/api_chlo.md.CTmmK6a-.lean.js
@@ -0,0 +1 @@
+import{_ as o,c as r,j as l,a as e,G as n,a2 as s,B as i,o as u}from"./chunks/framework.2yyKLD8d.js";const Rt=JSON.parse('{"title":"CHLO Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/chlo.md","filePath":"api/chlo.md","lastUpdated":null}'),d={name:"api/chlo.md"},R={class:"jldocstring custom-block"},p={class:"jldocstring custom-block"},c={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"},I={class:"jldocstring custom-block"},L={class:"jldocstring custom-block"},M={class:"jldocstring custom-block"},j={class:"jldocstring custom-block"},g={class:"jldocstring custom-block"},y={class:"jldocstring custom-block"},D={class:"jldocstring custom-block"},w={class:"jldocstring custom-block"},x={class:"jldocstring custom-block"},T={class:"jldocstring custom-block"},h={class:"jldocstring custom-block"},V={class:"jldocstring custom-block"},k={class:"jldocstring custom-block"},O={class:"jldocstring custom-block"},z={class:"jldocstring custom-block"},C={class:"jldocstring custom-block"},A={class:"jldocstring custom-block"},E={class:"jldocstring custom-block"},H={class:"jldocstring custom-block"},v={class:"jldocstring custom-block"},$={class:"jldocstring custom-block"},S={class:"jldocstring custom-block"},q={class:"jldocstring custom-block"},N={class:"jldocstring custom-block"},B={class:"jldocstring custom-block"},P={class:"jldocstring custom-block"},G={class:"jldocstring custom-block"},F={class:"jldocstring custom-block"},U={class:"jldocstring custom-block"},Z={class:"jldocstring custom-block"},J={class:"jldocstring custom-block"},K={class:"jldocstring custom-block"},Q={class:"jldocstring custom-block"},W={class:"jldocstring custom-block"},X={class:"jldocstring custom-block"},Y={class:"jldocstring custom-block"},_={class:"jldocstring custom-block"},tt={class:"jldocstring custom-block"},lt={class:"jldocstring custom-block"},et={class:"jldocstring custom-block"},at={class:"jldocstring custom-block"};function nt(st,t,ot,rt,it,ut){const a=i("Badge");return u(),r("div",null,[t[277]||(t[277]=l("h1",{id:"CHLO-Dialect",tabindex:"-1"},[e("CHLO Dialect "),l("a",{class:"header-anchor",href:"#CHLO-Dialect","aria-label":'Permalink to "CHLO Dialect {#CHLO-Dialect}"'},"")],-1)),t[278]||(t[278]=l("p",null,[e("Refer to the "),l("a",{href:"https://github.com/openxla/xla/tree/main/xla/mlir_hlo#hlo-client-dialect-chlo",target:"_blank",rel:"noreferrer"},"official documentation"),e(" for more details.")],-1)),l("details",R,[l("summary",null,[t[0]||(t[0]=l("a",{id:"Reactant.MLIR.Dialects.chlo._asin_acos_kernel-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo._asin_acos_kernel-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo._asin_acos_kernel")],-1)),t[1]||(t[1]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[2]||(t[2]=l("p",null,[l("code",null,"_asin_acos_kernel")],-1)),t[3]||(t[3]=l("p",null,[e("Returns "),l("code",null,"AsinAcosKernel(operand)"),e(" element-wise.")],-1)),t[4]||(t[4]=l("p",null,"If w = _asin_acos_kernel(z) w' = _asin_acos_kernel(I * z) then asin(z) = complex(atan2(z.real, w.real), sign(z.imag) * w.imag) acos(z) = complex(atan2(w.real, z.real), -sign(z.imag) * w.imag) asinh(z) = complex(sign(z.real) * w'.imag, atan2(z.imag, w'.real)) acosh(z) = complex(w.imag, sign(z.imag) * atan2(w.real, z.real))",-1)),t[5]||(t[5]=l("p",null,"This op is used as an intermediate value in decompositions and should never be constructed directly by frameworks or consumed by backends.",-1)),t[6]||(t[6]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L76-L93",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",p,[l("summary",null,[t[7]||(t[7]=l("a",{id:"Reactant.MLIR.Dialects.chlo.acos-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.acos-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.acos")],-1)),t[8]||(t[8]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[9]||(t[9]=l("p",null,[l("code",null,"acos")],-1)),t[10]||(t[10]=l("p",null,[e("Returns "),l("code",null,"Acos(operand)"),e(" element-wise.")],-1)),t[11]||(t[11]=l("p",null,"$",-1)),t[12]||(t[12]=l("p",null,"\\acos(x) = 2 * \\atan(\\sqrt(1 - x^2) / (1 + x)) if x != -1 = pi if x == -1 $",-1)),t[13]||(t[13]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L16-L25",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",c,[l("summary",null,[t[14]||(t[14]=l("a",{id:"Reactant.MLIR.Dialects.chlo.acosh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.acosh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.acosh")],-1)),t[15]||(t[15]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[16]||(t[16]=l("p",null,[l("code",null,"acosh")],-1)),t[17]||(t[17]=l("p",null,[e("Returns "),l("code",null,"Acosh(operand)"),e(" element-wise.")],-1)),t[18]||(t[18]=l("p",null,"$",-1)),t[19]||(t[19]=l("p",null,"\\acosh(x) = log(x + sqrt(x^2 - 1)) if x >= -1 \\acosh(x) = nan if x < -1 $",-1)),t[20]||(t[20]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L46-L55",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",m,[l("summary",null,[t[21]||(t[21]=l("a",{id:"Reactant.MLIR.Dialects.chlo.asin-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.asin-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.asin")],-1)),t[22]||(t[22]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[23]||(t[23]=l("p",null,[l("code",null,"asin")],-1)),t[24]||(t[24]=l("p",null,[e("Returns "),l("code",null,"Asin(operand)"),e(" element-wise.")],-1)),t[25]||(t[25]=l("p",null,"$",-1)),t[26]||(t[26]=l("p",null,"\\asin(x) = 2 * atan(x / (1 + sqrt(1 - x^2))) $",-1)),t[27]||(t[27]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L116-L124",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",f,[l("summary",null,[t[28]||(t[28]=l("a",{id:"Reactant.MLIR.Dialects.chlo.asinh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.asinh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.asinh")],-1)),t[29]||(t[29]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[30]||(t[30]=l("p",null,[l("code",null,"asinh")],-1)),t[31]||(t[31]=l("p",null,[e("Returns "),l("code",null,"Asinh(operand)"),e(" element-wise.")],-1)),t[32]||(t[32]=l("p",null,"$",-1)),t[33]||(t[33]=l("p",null,"\\asinh(x) = log(x + sqrt(x^2 + 1)) $",-1)),t[34]||(t[34]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L145-L153",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",b,[l("summary",null,[t[35]||(t[35]=l("a",{id:"Reactant.MLIR.Dialects.chlo.atan-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.atan-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.atan")],-1)),t[36]||(t[36]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[37]||(t[37]=l("p",null,[l("code",null,"atan")],-1)),t[38]||(t[38]=l("p",null,[e("Returns "),l("code",null,"Atan(operand)"),e(" element-wise.")],-1)),t[39]||(t[39]=l("p",null,"$",-1)),t[40]||(t[40]=l("p",null,"\\atan(x) = \\atan2(x, 1) $",-1)),t[41]||(t[41]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L174-L182",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",I,[l("summary",null,[t[42]||(t[42]=l("a",{id:"Reactant.MLIR.Dialects.chlo.atanh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.atanh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.atanh")],-1)),t[43]||(t[43]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[44]||(t[44]=l("p",null,[l("code",null,"atanh")],-1)),t[45]||(t[45]=l("p",null,[e("Returns "),l("code",null,"Atanh(operand)"),e(" element-wise.")],-1)),t[46]||(t[46]=l("p",null,"$",-1)),t[47]||(t[47]=l("p",null,"\\atanh(x) = 0.5 * log((1 + x) / (1 - x)) if abs(x) <= 1 = nan otherwise $",-1)),t[48]||(t[48]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L203-L212",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",L,[l("summary",null,[t[49]||(t[49]=l("a",{id:"Reactant.MLIR.Dialects.chlo.bessel_i1e-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.bessel_i1e-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.bessel_i1e")],-1)),t[50]||(t[50]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[51]||(t[51]=l("p",null,[l("code",null,"bessel_i1e")],-1)),t[52]||(t[52]=l("p",null,[e("Returns "),l("code",null,"bessel_i1e(operand)"),e(" element-wise.")],-1)),t[53]||(t[53]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L233-L237",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",M,[l("summary",null,[t[54]||(t[54]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_add-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_add-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_add")],-1)),t[55]||(t[55]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[56]||(t[56]=l("p",null,[l("code",null,"broadcast_add")],-1)),t[57]||(t[57]=l("p",null,[e("Returns "),l("code",null,"lhs + rhs"),e(" element-wise.")],-1)),t[58]||(t[58]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[59]||(t[59]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L260-L267",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",j,[l("summary",null,[t[60]||(t[60]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_and-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_and-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_and")],-1)),t[61]||(t[61]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[62]||(t[62]=l("p",null,[l("code",null,"broadcast_and")],-1)),t[63]||(t[63]=l("p",null,[e("Returns "),l("code",null,"logical_and(lhs, rhs)"),e(" element-wise.")],-1)),t[64]||(t[64]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[65]||(t[65]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L296-L303",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",g,[l("summary",null,[t[66]||(t[66]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_atan2-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_atan2-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_atan2")],-1)),t[67]||(t[67]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[68]||(t[68]=l("p",null,[l("code",null,"broadcast_atan2")],-1)),t[69]||(t[69]=l("p",null,[e("Returns "),l("code",null,"atan2(lhs/rhs)"),e(" element-wise.")],-1)),t[70]||(t[70]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[71]||(t[71]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L332-L339",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",y,[l("summary",null,[t[72]||(t[72]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_compare-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_compare-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_compare")],-1)),t[73]||(t[73]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[74]||(t[74]=s('
broadcast_compare
Compares lhs and rhs elementwise according to comparison_direction and compare_type. If unspecified, compare_type is FLOAT for float element types, SIGNED for signed element types and UNSIGNED for unsigned element types.
',4))]),l("details",D,[l("summary",null,[t[75]||(t[75]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_complex-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_complex-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_complex")],-1)),t[76]||(t[76]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[77]||(t[77]=l("p",null,[l("code",null,"broadcast_complex")],-1)),t[78]||(t[78]=l("p",null,"Performs element-wise conversion of a pair of real and imaginary values to a complex value.",-1)),t[79]||(t[79]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L413-L418",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",w,[l("summary",null,[t[80]||(t[80]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_divide-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_divide-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_divide")],-1)),t[81]||(t[81]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[82]||(t[82]=l("p",null,[l("code",null,"broadcast_divide")],-1)),t[83]||(t[83]=l("p",null,[e("Returns "),l("code",null,"lhs / rhs"),e(" element-wise.")],-1)),t[84]||(t[84]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[85]||(t[85]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L447-L454",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",x,[l("summary",null,[t[86]||(t[86]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_maximum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_maximum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_maximum")],-1)),t[87]||(t[87]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[88]||(t[88]=l("p",null,[l("code",null,"broadcast_maximum")],-1)),t[89]||(t[89]=l("p",null,[e("Returns "),l("code",null,"max(lhs, rhs)"),e(" element-wise.")],-1)),t[90]||(t[90]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[91]||(t[91]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L483-L490",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",T,[l("summary",null,[t[92]||(t[92]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_minimum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_minimum-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_minimum")],-1)),t[93]||(t[93]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[94]||(t[94]=l("p",null,[l("code",null,"broadcast_minimum")],-1)),t[95]||(t[95]=l("p",null,[e("Returns "),l("code",null,"min(lhs, rhs)"),e(" element-wise.")],-1)),t[96]||(t[96]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[97]||(t[97]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L519-L526",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",h,[l("summary",null,[t[98]||(t[98]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_multiply-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_multiply-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_multiply")],-1)),t[99]||(t[99]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[100]||(t[100]=l("p",null,[l("code",null,"broadcast_multiply")],-1)),t[101]||(t[101]=l("p",null,[e("Returns "),l("code",null,"lhs * rhs"),e(" element-wise.")],-1)),t[102]||(t[102]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[103]||(t[103]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L555-L562",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",V,[l("summary",null,[t[104]||(t[104]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_next_after")],-1)),t[105]||(t[105]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[106]||(t[106]=l("p",null,[l("code",null,"broadcast_next_after")],-1)),t[107]||(t[107]=l("p",null,[e("Returns the next representable value of "),l("code",null,"lhs"),e(" in the direction of "),l("code",null,"rhs"),e(", element-wise. It can also return a subnormal number.")],-1)),t[108]||(t[108]=l("p",null,"Equivalent to the C++ std::nextafter function.",-1)),t[109]||(t[109]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L591-L598",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",k,[l("summary",null,[t[110]||(t[110]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_or-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_or-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_or")],-1)),t[111]||(t[111]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[112]||(t[112]=l("p",null,[l("code",null,"broadcast_or")],-1)),t[113]||(t[113]=l("p",null,[e("Returns "),l("code",null,"logical_or(lhs, rhs)"),e(" element-wise.")],-1)),t[114]||(t[114]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[115]||(t[115]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L627-L634",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",O,[l("summary",null,[t[116]||(t[116]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_polygamma")],-1)),t[117]||(t[117]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[118]||(t[118]=l("p",null,[l("code",null,"broadcast_polygamma")],-1)),t[119]||(t[119]=l("p",null,[e("Returns "),l("code",null,"Polygamma(operand, operand)"),e(" element-wise.")],-1)),t[120]||(t[120]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L663-L667",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",z,[l("summary",null,[t[121]||(t[121]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_power-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_power-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_power")],-1)),t[122]||(t[122]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[123]||(t[123]=l("p",null,[l("code",null,"broadcast_power")],-1)),t[124]||(t[124]=l("p",null,[e("Returns "),l("code",null,"lhs ^ rhs"),e(" element-wise.")],-1)),t[125]||(t[125]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[126]||(t[126]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L696-L703",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",C,[l("summary",null,[t[127]||(t[127]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_remainder-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_remainder-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_remainder")],-1)),t[128]||(t[128]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[129]||(t[129]=l("p",null,[l("code",null,"broadcast_remainder")],-1)),t[130]||(t[130]=l("p",null,[e("Returns "),l("code",null,"lhs % rhs"),e(" element-wise.")],-1)),t[131]||(t[131]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[132]||(t[132]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L732-L739",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",A,[l("summary",null,[t[133]||(t[133]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_select-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_select-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_select")],-1)),t[134]||(t[134]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[135]||(t[135]=l("p",null,[l("code",null,"broadcast_select")],-1)),t[136]||(t[136]=l("p",null,"Constructs an output array from elements of two input arrays, based on the values of a predicate array.",-1)),t[137]||(t[137]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#select",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#select")],-1)),t[138]||(t[138]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L768-L775",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",E,[l("summary",null,[t[139]||(t[139]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_shift_left-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_shift_left-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_shift_left")],-1)),t[140]||(t[140]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[141]||(t[141]=l("p",null,[l("code",null,"broadcast_shift_left")],-1)),t[142]||(t[142]=l("p",null,[e("Returns "),l("code",null,"lhs << rhs"),e(" element-wise.")],-1)),t[143]||(t[143]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[144]||(t[144]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L802-L809",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",H,[l("summary",null,[t[145]||(t[145]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_arithmetic-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_shift_right_arithmetic-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_arithmetic")],-1)),t[146]||(t[146]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[147]||(t[147]=l("p",null,[l("code",null,"broadcast_shift_right_arithmetic")],-1)),t[148]||(t[148]=l("p",null,[e("Returns "),l("code",null,"lhs >> rhs"),e(" element-wise.")],-1)),t[149]||(t[149]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[150]||(t[150]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L838-L845",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",v,[l("summary",null,[t[151]||(t[151]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_logical-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_shift_right_logical-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_shift_right_logical")],-1)),t[152]||(t[152]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[153]||(t[153]=l("p",null,[l("code",null,"broadcast_shift_right_logical")],-1)),t[154]||(t[154]=l("p",null,[e("Returns "),l("code",null,"lhs >> rhs"),e(" element-wise.")],-1)),t[155]||(t[155]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[156]||(t[156]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L874-L881",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",$,[l("summary",null,[t[157]||(t[157]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_subtract-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_subtract-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_subtract")],-1)),t[158]||(t[158]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[159]||(t[159]=l("p",null,[l("code",null,"broadcast_subtract")],-1)),t[160]||(t[160]=l("p",null,[e("Returns "),l("code",null,"lhs - rhs"),e(" element-wise.")],-1)),t[161]||(t[161]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[162]||(t[162]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L910-L917",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",S,[l("summary",null,[t[163]||(t[163]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_xor-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_xor-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_xor")],-1)),t[164]||(t[164]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[165]||(t[165]=l("p",null,[l("code",null,"broadcast_xor")],-1)),t[166]||(t[166]=l("p",null,[e("Returns "),l("code",null,"logical_xor(lhs, rhs)"),e(" element-wise.")],-1)),t[167]||(t[167]=l("p",null,[e("See "),l("a",{href:"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations",target:"_blank",rel:"noreferrer"},"https://www.tensorflow.org/xla/operation_semantics#element-wise_binary_arithmetic_operations"),e(".")],-1)),t[168]||(t[168]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L946-L953",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",q,[l("summary",null,[t[169]||(t[169]=l("a",{id:"Reactant.MLIR.Dialects.chlo.broadcast_zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.broadcast_zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.broadcast_zeta")],-1)),t[170]||(t[170]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[171]||(t[171]=l("p",null,[l("code",null,"broadcast_zeta")],-1)),t[172]||(t[172]=l("p",null,[e("Returns "),l("code",null,"Zeta(operand, operand)"),e(" element-wise.")],-1)),t[173]||(t[173]=l("p",null,"$",-1)),t[174]||(t[174]=l("p",null,"(\\zeta(x, q) = \\sum_{n=0}^{\\infty} (q + n)^{-x}) $",-1)),t[175]||(t[175]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L982-L990",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",N,[l("summary",null,[t[176]||(t[176]=l("a",{id:"Reactant.MLIR.Dialects.chlo.conj-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.conj-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.conj")],-1)),t[177]||(t[177]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[178]||(t[178]=l("p",null,[l("code",null,"conj")],-1)),t[179]||(t[179]=l("p",null,[e("Returns "),l("code",null,"Conj(operand)"),e(" element-wise.")],-1)),t[180]||(t[180]=l("p",null,"$",-1)),t[181]||(t[181]=l("p",null,"\\conj(x) = (\\real(x), \\neg(\\imag(x))) $",-1)),t[182]||(t[182]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1019-L1027",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",B,[l("summary",null,[t[183]||(t[183]=l("a",{id:"Reactant.MLIR.Dialects.chlo.constant-Tuple{}",href:"#Reactant.MLIR.Dialects.chlo.constant-Tuple{}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.constant")],-1)),t[184]||(t[184]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[185]||(t[185]=l("p",null,[l("code",null,"constant")],-1)),t[186]||(t[186]=l("p",null,"Represents a constant value.",-1)),t[187]||(t[187]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1075-L1079",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",P,[l("summary",null,[t[188]||(t[188]=l("a",{id:"Reactant.MLIR.Dialects.chlo.constant_like-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.constant_like-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.constant_like")],-1)),t[189]||(t[189]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[190]||(t[190]=l("p",null,[l("code",null,"constant_like")],-1)),t[191]||(t[191]=l("p",null,"Returns a splat constant of the same shape as the operand.",-1)),t[192]||(t[192]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1048-L1052",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",G,[l("summary",null,[t[193]||(t[193]=l("a",{id:"Reactant.MLIR.Dialects.chlo.cosh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.cosh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.cosh")],-1)),t[194]||(t[194]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[195]||(t[195]=l("p",null,[l("code",null,"cosh")],-1)),t[196]||(t[196]=l("p",null,[e("Returns "),l("code",null,"Cosh(operand)"),e(" element-wise.")],-1)),t[197]||(t[197]=l("p",null,"$",-1)),t[198]||(t[198]=l("p",null,"\\cosh(x) = (e^x + e^-x) / 2 $",-1)),t[199]||(t[199]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1100-L1108",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",F,[l("summary",null,[t[200]||(t[200]=l("a",{id:"Reactant.MLIR.Dialects.chlo.digamma-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.digamma-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.digamma")],-1)),t[201]||(t[201]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[202]||(t[202]=l("p",null,[l("code",null,"digamma")],-1)),t[203]||(t[203]=l("p",null,[e("Returns "),l("code",null,"Digamma(operand)"),e(" element-wise.")],-1)),t[204]||(t[204]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1129-L1133",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",U,[l("summary",null,[t[205]||(t[205]=l("a",{id:"Reactant.MLIR.Dialects.chlo.erf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.erf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.erf")],-1)),t[206]||(t[206]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[207]||(t[207]=l("p",null,[l("code",null,"erf")],-1)),t[208]||(t[208]=l("p",null,[e("Computes the Gauss error function of "),l("code",null,"x"),e(" element-wise.")],-1)),t[209]||(t[209]=l("p",null,"erf(x) = erf_impl(x) if |x| < 1 = 1 - erfc_impl(x) otherwise",-1)),t[210]||(t[210]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1183-L1190",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",Z,[l("summary",null,[t[211]||(t[211]=l("a",{id:"Reactant.MLIR.Dialects.chlo.erf_inv-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.erf_inv-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.erf_inv")],-1)),t[212]||(t[212]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[213]||(t[213]=l("p",null,[l("code",null,"erf_inv")],-1)),t[214]||(t[214]=l("p",null,[e("Returns "),l("code",null,"ErfInv(operand)"),e(" element-wise.")],-1)),t[215]||(t[215]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1156-L1160",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",J,[l("summary",null,[t[216]||(t[216]=l("a",{id:"Reactant.MLIR.Dialects.chlo.erfc-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.erfc-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.erfc")],-1)),t[217]||(t[217]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[218]||(t[218]=l("p",null,[l("code",null,"erfc")],-1)),t[219]||(t[219]=l("p",null,"Computes an approximation of the error function complement (1 - erf(x)).",-1)),t[220]||(t[220]=l("p",null,"erfc(x) = erfc_impl(x) if |x| > 1 = 1 - erf_impl(x) otherwise",-1)),t[221]||(t[221]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1211-L1218",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",K,[l("summary",null,[t[222]||(t[222]=l("a",{id:"Reactant.MLIR.Dialects.chlo.is_inf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.is_inf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.is_inf")],-1)),t[223]||(t[223]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[224]||(t[224]=l("p",null,[l("code",null,"is_inf")],-1)),t[225]||(t[225]=l("p",null,"Returns if a value is +/-inf element-wise.",-1)),t[226]||(t[226]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1239-L1243",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",Q,[l("summary",null,[t[227]||(t[227]=l("a",{id:"Reactant.MLIR.Dialects.chlo.is_neg_inf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.is_neg_inf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.is_neg_inf")],-1)),t[228]||(t[228]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[229]||(t[229]=l("p",null,[l("code",null,"is_neg_inf")],-1)),t[230]||(t[230]=l("p",null,"Returns if a value is -inf element-wise.",-1)),t[231]||(t[231]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1264-L1268",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",W,[l("summary",null,[t[232]||(t[232]=l("a",{id:"Reactant.MLIR.Dialects.chlo.is_pos_inf-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.is_pos_inf-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.is_pos_inf")],-1)),t[233]||(t[233]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[234]||(t[234]=l("p",null,[l("code",null,"is_pos_inf")],-1)),t[235]||(t[235]=l("p",null,"Returns if a value is +inf element-wise.",-1)),t[236]||(t[236]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1291-L1295",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",X,[l("summary",null,[t[237]||(t[237]=l("a",{id:"Reactant.MLIR.Dialects.chlo.lgamma-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.lgamma-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.lgamma")],-1)),t[238]||(t[238]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[239]||(t[239]=l("p",null,[l("code",null,"lgamma")],-1)),t[240]||(t[240]=l("p",null,[e("Returns "),l("code",null,"Lgamma(operand)"),e(" element-wise.")],-1)),t[241]||(t[241]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1318-L1322",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",Y,[l("summary",null,[t[242]||(t[242]=l("a",{id:"Reactant.MLIR.Dialects.chlo.next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.next_after-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.next_after")],-1)),t[243]||(t[243]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[244]||(t[244]=l("p",null,[l("code",null,"next_after")],-1)),t[245]||(t[245]=l("p",null,[e("Returns the next representable value of "),l("code",null,"x"),e(" in the direction of "),l("code",null,"y"),e(", element-wise. It can also return a subnormal number.")],-1)),t[246]||(t[246]=l("p",null,"Equivalent to the C++ std::nextafter function.",-1)),t[247]||(t[247]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1343-L1350",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",_,[l("summary",null,[t[248]||(t[248]=l("a",{id:"Reactant.MLIR.Dialects.chlo.polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.polygamma-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.polygamma")],-1)),t[249]||(t[249]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[250]||(t[250]=l("p",null,[l("code",null,"polygamma")],-1)),t[251]||(t[251]=l("p",null,[e("Returns "),l("code",null,"Polygamma(operand, operand)"),e(" element-wise.")],-1)),t[252]||(t[252]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1373-L1377",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",tt,[l("summary",null,[t[253]||(t[253]=l("a",{id:"Reactant.MLIR.Dialects.chlo.sinh-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.sinh-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.sinh")],-1)),t[254]||(t[254]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[255]||(t[255]=l("p",null,[l("code",null,"sinh")],-1)),t[256]||(t[256]=l("p",null,[e("Returns "),l("code",null,"Sinh(operand)"),e(" element-wise.")],-1)),t[257]||(t[257]=l("p",null,"$",-1)),t[258]||(t[258]=l("p",null,"\\sinh(x) = (e^x - e^-x) / 2 if |x| < 1 = e^(x + log(1/2)) - e^(-x + log(1/2)) otherwise. $",-1)),t[259]||(t[259]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1400-L1409",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",lt,[l("summary",null,[t[260]||(t[260]=l("a",{id:"Reactant.MLIR.Dialects.chlo.tan-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.tan-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.tan")],-1)),t[261]||(t[261]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[262]||(t[262]=l("p",null,[l("code",null,"tan")],-1)),t[263]||(t[263]=l("p",null,[e("Returns "),l("code",null,"Tan(operand)"),e(" element-wise.")],-1)),t[264]||(t[264]=l("p",null,"$",-1)),t[265]||(t[265]=l("p",null,"\\tan(x) = \\sin(x) / \\cos(x) $",-1)),t[266]||(t[266]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1430-L1438",target:"_blank",rel:"noreferrer"},"source")],-1))]),l("details",et,[l("summary",null,[t[267]||(t[267]=l("a",{id:"Reactant.MLIR.Dialects.chlo.top_k-Tuple{Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.top_k-Tuple{Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.top_k")],-1)),t[268]||(t[268]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[269]||(t[269]=s('
top_k
If the input is a vector (rank-1), finds the k largest entries in the vector and outputs their values and indices as vectors. Thus values[j] is the j-th largest entry in input, and its index is indices[j].
For matrices (resp. higher rank input), computes the top k entries in each row (resp. vector along the last dimension). Thus,
',6))]),l("details",at,[l("summary",null,[t[270]||(t[270]=l("a",{id:"Reactant.MLIR.Dialects.chlo.zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}",href:"#Reactant.MLIR.Dialects.chlo.zeta-Tuple{Reactant.MLIR.IR.Value, Reactant.MLIR.IR.Value}"},[l("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.chlo.zeta")],-1)),t[271]||(t[271]=e()),n(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[272]||(t[272]=l("p",null,[l("code",null,"zeta")],-1)),t[273]||(t[273]=l("p",null,[e("Returns "),l("code",null,"Zeta(operand, operand)"),e(" element-wise.")],-1)),t[274]||(t[274]=l("p",null,"$",-1)),t[275]||(t[275]=l("p",null,"(\\zeta(x, q) = \\sum_{n=0}^{\\infty} (q + n)^{-x}) $",-1)),t[276]||(t[276]=l("p",null,[l("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/CHLO.jl#L1500-L1508",target:"_blank",rel:"noreferrer"},"source")],-1))])])}const pt=o(d,[["render",nt]]);export{Rt as __pageData,pt as default};
diff --git a/previews/PR363/assets/api_enzyme.md.C8cSUFJK.js b/previews/PR363/assets/api_enzyme.md.C8cSUFJK.js
new file mode 100644
index 00000000..b4cc968c
--- /dev/null
+++ b/previews/PR363/assets/api_enzyme.md.C8cSUFJK.js
@@ -0,0 +1 @@
+import{_ as l,c as o,j as t,a,G as s,B as r,o as d}from"./chunks/framework.2yyKLD8d.js";const R=JSON.parse('{"title":"Enzyme Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/enzyme.md","filePath":"api/enzyme.md","lastUpdated":null}'),i={name:"api/enzyme.md"},c={class:"jldocstring custom-block"};function m(p,e,u,y,f,z){const n=r("Badge");return d(),o("div",null,[e[5]||(e[5]=t("h1",{id:"Enzyme-Dialect",tabindex:"-1"},[a("Enzyme Dialect "),t("a",{class:"header-anchor",href:"#Enzyme-Dialect","aria-label":'Permalink to "Enzyme Dialect {#Enzyme-Dialect}"'},"")],-1)),t("details",c,[t("summary",null,[e[0]||(e[0]=t("a",{id:"Reactant.MLIR.Dialects.enzyme.addTo-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.enzyme.addTo-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[t("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.enzyme.addTo")],-1)),e[1]||(e[1]=a()),s(n,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=t("p",null,[t("code",null,"addTo")],-1)),e[3]||(e[3]=t("p",null,"TODO",-1)),e[4]||(e[4]=t("p",null,[t("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Enzyme.jl#L16-L20",target:"_blank",rel:"noreferrer"},"source")],-1))])])}const b=l(i,[["render",m]]);export{R as __pageData,b as default};
diff --git a/previews/PR363/assets/api_enzyme.md.C8cSUFJK.lean.js b/previews/PR363/assets/api_enzyme.md.C8cSUFJK.lean.js
new file mode 100644
index 00000000..b4cc968c
--- /dev/null
+++ b/previews/PR363/assets/api_enzyme.md.C8cSUFJK.lean.js
@@ -0,0 +1 @@
+import{_ as l,c as o,j as t,a,G as s,B as r,o as d}from"./chunks/framework.2yyKLD8d.js";const R=JSON.parse('{"title":"Enzyme Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/enzyme.md","filePath":"api/enzyme.md","lastUpdated":null}'),i={name:"api/enzyme.md"},c={class:"jldocstring custom-block"};function m(p,e,u,y,f,z){const n=r("Badge");return d(),o("div",null,[e[5]||(e[5]=t("h1",{id:"Enzyme-Dialect",tabindex:"-1"},[a("Enzyme Dialect "),t("a",{class:"header-anchor",href:"#Enzyme-Dialect","aria-label":'Permalink to "Enzyme Dialect {#Enzyme-Dialect}"'},"")],-1)),t("details",c,[t("summary",null,[e[0]||(e[0]=t("a",{id:"Reactant.MLIR.Dialects.enzyme.addTo-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.enzyme.addTo-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[t("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.enzyme.addTo")],-1)),e[1]||(e[1]=a()),s(n,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),e[2]||(e[2]=t("p",null,[t("code",null,"addTo")],-1)),e[3]||(e[3]=t("p",null,"TODO",-1)),e[4]||(e[4]=t("p",null,[t("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/Dialects/Enzyme.jl#L16-L20",target:"_blank",rel:"noreferrer"},"source")],-1))])])}const b=l(i,[["render",m]]);export{R as __pageData,b as default};
diff --git a/previews/PR363/assets/api_func.md.BwGv464S.js b/previews/PR363/assets/api_func.md.BwGv464S.js
new file mode 100644
index 00000000..5d05fa50
--- /dev/null
+++ b/previews/PR363/assets/api_func.md.BwGv464S.js
@@ -0,0 +1,26 @@
+import{_ as c,c as i,j as e,a,G as s,a2 as l,B as o,o as p}from"./chunks/framework.2yyKLD8d.js";const L=JSON.parse('{"title":"Func Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/func.md","filePath":"api/func.md","lastUpdated":null}'),r={name:"api/func.md"},u={class:"jldocstring custom-block"},d={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"};function g(h,n,R,y,v,x){const t=o("Badge");return p(),i("div",null,[n[15]||(n[15]=e("h1",{id:"Func-Dialect",tabindex:"-1"},[a("Func Dialect "),e("a",{class:"header-anchor",href:"#Func-Dialect","aria-label":'Permalink to "Func Dialect {#Func-Dialect}"'},"")],-1)),n[16]||(n[16]=e("p",null,[a("Refer to the "),e("a",{href:"https://mlir.llvm.org/docs/Dialects/Func/",target:"_blank",rel:"noreferrer"},"official documentation"),a(" for more details.")],-1)),e("details",u,[e("summary",null,[n[0]||(n[0]=e("a",{id:"Reactant.MLIR.Dialects.func.call-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.func.call-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[e("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.func.call")],-1)),n[1]||(n[1]=a()),s(t,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),n[2]||(n[2]=l('
call
The func.call operation represents a direct call to a function that is within the same symbol scope as the call. The operands and result types of the call must match the specified function type. The callee is encoded as a symbol reference attribute named "callee".
The func.call_indirect operation represents an indirect call to a value of function type. The operands and result types of the call must match the specified function type.
The func.constant operation produces an SSA value from a symbol reference to a func.func operation
Example
mlir
// Reference to function @myfn.
+%2 = func.constant @myfn : (tensor<16xf32>, f32) -> tensor<16xf32>
+
+// Equivalent generic forms
+%2 = "func.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).
Operations within the function cannot implicitly capture values defined outside of the function, i.e. Functions are IsolatedFromAbove. All external references must use function arguments or attributes that establish a symbolic connection (e.g. symbols referenced by name via a string attribute like SymbolRefAttr). An external function declaration (used when referring to a function declared in some other module) has no body. While the MLIR textual form provides a nice inline syntax for function arguments, they are internally represented as “block arguments” to the first block in the region.
Only dialect attribute names may be specified in the attribute dictionaries for function arguments, results, or the function itself.
Example
mlir
// External function definitions.
+func.func private @abort()
+func.func private @scribble(i32, i64, memref<? x 128 x f32, #layout_map0>) -> f64
+
+// A function that returns its argument twice:
+func.func @count(%x: i64) -> (i64, i64)
+ attributes {fruit = "banana"} {
+ return %x, %x: i64, i64
+}
+
+// A function with an argument attribute
+func.func private @example_fn_arg(%x: i32 {swift.self = unit})
+
+// A function with a result attribute
+func.func private @example_fn_result() -> (f64 {dialectName.attrName = 0 : i64})
+
+// A function with an attribute
+func.func private @example_fn_attr() attributes {dialectName.attrName = false}
The func.return operation represents a return operation within a function. The operation takes variable number of operands and produces no results. The operand number and types must match the signature of the function that contains the operation.
`,5))])])}const M=c(r,[["render",g]]);export{L as __pageData,M as default};
diff --git a/previews/PR363/assets/api_func.md.BwGv464S.lean.js b/previews/PR363/assets/api_func.md.BwGv464S.lean.js
new file mode 100644
index 00000000..5d05fa50
--- /dev/null
+++ b/previews/PR363/assets/api_func.md.BwGv464S.lean.js
@@ -0,0 +1,26 @@
+import{_ as c,c as i,j as e,a,G as s,a2 as l,B as o,o as p}from"./chunks/framework.2yyKLD8d.js";const L=JSON.parse('{"title":"Func Dialect","description":"","frontmatter":{},"headers":[],"relativePath":"api/func.md","filePath":"api/func.md","lastUpdated":null}'),r={name:"api/func.md"},u={class:"jldocstring custom-block"},d={class:"jldocstring custom-block"},f={class:"jldocstring custom-block"},m={class:"jldocstring custom-block"},b={class:"jldocstring custom-block"};function g(h,n,R,y,v,x){const t=o("Badge");return p(),i("div",null,[n[15]||(n[15]=e("h1",{id:"Func-Dialect",tabindex:"-1"},[a("Func Dialect "),e("a",{class:"header-anchor",href:"#Func-Dialect","aria-label":'Permalink to "Func Dialect {#Func-Dialect}"'},"")],-1)),n[16]||(n[16]=e("p",null,[a("Refer to the "),e("a",{href:"https://mlir.llvm.org/docs/Dialects/Func/",target:"_blank",rel:"noreferrer"},"official documentation"),a(" for more details.")],-1)),e("details",u,[e("summary",null,[n[0]||(n[0]=e("a",{id:"Reactant.MLIR.Dialects.func.call-Tuple{Vector{Reactant.MLIR.IR.Value}}",href:"#Reactant.MLIR.Dialects.func.call-Tuple{Vector{Reactant.MLIR.IR.Value}}"},[e("span",{class:"jlbinding"},"Reactant.MLIR.Dialects.func.call")],-1)),n[1]||(n[1]=a()),s(t,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),n[2]||(n[2]=l('
call
The func.call operation represents a direct call to a function that is within the same symbol scope as the call. The operands and result types of the call must match the specified function type. The callee is encoded as a symbol reference attribute named "callee".
The func.call_indirect operation represents an indirect call to a value of function type. The operands and result types of the call must match the specified function type.
The func.constant operation produces an SSA value from a symbol reference to a func.func operation
Example
mlir
// Reference to function @myfn.
+%2 = func.constant @myfn : (tensor<16xf32>, f32) -> tensor<16xf32>
+
+// Equivalent generic forms
+%2 = "func.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).
Operations within the function cannot implicitly capture values defined outside of the function, i.e. Functions are IsolatedFromAbove. All external references must use function arguments or attributes that establish a symbolic connection (e.g. symbols referenced by name via a string attribute like SymbolRefAttr). An external function declaration (used when referring to a function declared in some other module) has no body. While the MLIR textual form provides a nice inline syntax for function arguments, they are internally represented as “block arguments” to the first block in the region.
Only dialect attribute names may be specified in the attribute dictionaries for function arguments, results, or the function itself.
Example
mlir
// External function definitions.
+func.func private @abort()
+func.func private @scribble(i32, i64, memref<? x 128 x f32, #layout_map0>) -> f64
+
+// A function that returns its argument twice:
+func.func @count(%x: i64) -> (i64, i64)
+ attributes {fruit = "banana"} {
+ return %x, %x: i64, i64
+}
+
+// A function with an argument attribute
+func.func private @example_fn_arg(%x: i32 {swift.self = unit})
+
+// A function with a result attribute
+func.func private @example_fn_result() -> (f64 {dialectName.attrName = 0 : i64})
+
+// A function with an attribute
+func.func private @example_fn_attr() attributes {dialectName.attrName = false}
The func.return operation represents a return operation within a function. The operation takes variable number of operands and produces no results. The operand number and types must match the signature of the function that contains the operation.
Creates an affine map with results defined by the given list of affine expressions. The map resulting map also has the requested number of input dimensions and symbols, regardless of them being used in the results.
',3))]),e("details",v,[e("summary",null,[t[48]||(t[48]=e("a",{id:"Reactant.MLIR.IR.Attribute-Tuple{T} where T<:AbstractFloat",href:"#Reactant.MLIR.IR.Attribute-Tuple{T} where T<:AbstractFloat"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Attribute")],-1)),t[49]||(t[49]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[50]||(t[50]=l('
Creates a floating point attribute in the given context with the given double value and double-precision FP semantics. If check=true, emits appropriate diagnostics on illegal arguments.
',3))]),e("details",T,[e("summary",null,[t[51]||(t[51]=e("a",{id:"Reactant.MLIR.IR.Attribute-Tuple{T} where T<:Complex",href:"#Reactant.MLIR.IR.Attribute-Tuple{T} where T<:Complex"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Attribute")],-1)),t[52]||(t[52]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[53]||(t[53]=l('
',3))]),e("details",F,[e("summary",null,[t[60]||(t[60]=e("a",{id:"Reactant.MLIR.IR.Attribute-Union{Tuple{T}, Tuple{T, Any}} where T<:Integer",href:"#Reactant.MLIR.IR.Attribute-Union{Tuple{T}, Tuple{T, Any}} where T<:Integer"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Attribute")],-1)),t[61]||(t[61]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[62]||(t[62]=l('
julia
Attribute(int)
Creates an integer attribute of the given type with the given integer value.
Creates an ExecutionEngine for the provided ModuleOp. The ModuleOp is expected to be "translatable" to LLVM IR (only contains operations in dialects that implement the LLVMTranslationDialectInterface). The module ownership stays with the client and can be destroyed as soon as the call returns. optLevel is the optimization level to be used for transformation and code generation. LLVM passes at optLevel are run before code generation. The number and array of paths corresponding to shared libraries that will be loaded are specified via numPaths and sharedLibPaths respectively. TODO: figure out other options.
Gets or creates a new integer set in the given context. The set is defined by a list of affine constraints, with the given number of input dimensions and symbols, which are treated as either equalities (eqflags is 1) or inequalities (eqflags is 0). Both constraints and eqflags need to be arrays of the same length.
A logical result value, essentially a boolean with named states. LLVM convention for using boolean values to designate success or failure of an operation is a moving target, so MLIR opted for an explicit class. Instances of LogicalResult must only be inspected using the associated functions.
Nest an OpPassManager under the provided OpPassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed.
Nest an OpPassManager under the top-level PassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed. To further nest more OpPassManager under the newly returned one, see mlirOpPassManagerNest below.
`,3))]),e("details",rt,[e("summary",null,[t[153]||(t[153]=e("a",{id:"Reactant.MLIR.IR.Type-Union{Tuple{Type{Complex{T}}}, Tuple{T}} where T",href:"#Reactant.MLIR.IR.Type-Union{Tuple{Type{Complex{T}}}, Tuple{T}} where T"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Type")],-1)),t[154]||(t[154]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[155]||(t[155]=l('
julia
Type(Complex{T}) where {T}
Creates a complex type with the given element type in the same context as the element type. The type is owned by the context.
Checks if two integer set objects are equal. This is a "shallow" comparison of two objects. Only the sets with some small number of constraints are uniqued and compare equal here. Set objects that represent the same integer set with different constraints may be considered non-equal by this check. Set difference followed by an (expensive) emptiness check should be used to check equivalence of the underlying integer sets.
Takes an operation owned by the caller and inserts it as index to the block. This is an expensive operation that scans the block linearly, prefer insertBefore/After instead.
Takes a block owned by the caller and inserts it at index to the given region. This is an expensive operation that linearly scans the region, prefer insertAfter/Before instead.
Inserts the given operation into the given symbol table. The operation must have the symbol trait. If the symbol table already has a symbol with the same name, renames the symbol being inserted to ensure name uniqueness. Note that this does not move the operation itself into the block of the symbol table operation, this should be done separately. Returns the name of the symbol after insertion.
Apply AffineExpr::replace(map) to each of the results and return a new new AffineMap with the new results and the specified number of dims and symbols.
Gets or creates a new integer set in which the values and dimensions of the given set are replaced with the given affine expressions. dimReplacements and symbolReplacements are expected to point to at least as many consecutive expressions as the given set has dimensions and symbols, respectively. The new set will have numResultDims and numResultSymbols dimensions and symbols, respectively.
Creates a dense elements attribute that has the same data as the given dense elements attribute and a different shaped type. The new type must have the same total number of elements.
',3))]),e("details",Jt,[e("summary",null,[t[273]||(t[273]=e("a",{id:"Reactant.MLIR.IR.BFloat16Type-Tuple{}",href:"#Reactant.MLIR.IR.BFloat16Type-Tuple{}"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.BFloat16Type")],-1)),t[274]||(t[274]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[275]||(t[275]=e("p",null,"BFloat16Type(; context=context())",-1)),t[276]||(t[276]=e("p",null,"Creates a bf16 type in the given context. The type is owned by the context.",-1)),t[277]||(t[277]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/IR/Type.jl#L157-L161",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Kt,[e("summary",null,[t[278]||(t[278]=e("a",{id:"Reactant.MLIR.IR.ConstantAffineMap-Tuple{Any}",href:"#Reactant.MLIR.IR.ConstantAffineMap-Tuple{Any}"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.ConstantAffineMap")],-1)),t[279]||(t[279]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[280]||(t[280]=l('
julia
ConstantAffineMap(val; context=context())
Creates a single constant result affine map in the context. The affine map is owned by the context.
Creates a dense elements attribute with the given shaped type from elements of a specific type. Expects the element type of the shaped type to match the data element type.
Creates a MemRef type with the given rank and shape, a potentially empty list of affine layout maps, the given memory space and element type, in the same context as element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates a MemRef type with the given rank, shape, memory space and element type in the same context as the element type. The type has no affine maps, i.e. represents a default row-major contiguous memref. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an Unranked MemRef type with the given element type and in the given memory space. The type is owned by the context of element type. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an identity affine map on the most minor dimensions in the context. The affine map is owned by the context. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an opaque attribute in the given context associated with the dialect identified by its namespace. The attribute contains opaque byte data of the specified length (data need not be null-terminated).
Creates an opaque type in the given context associated with the dialect identified by its namespace. The type contains opaque byte data of the specified length (data need not be null-terminated).
Creates an affine map with a permutation expression and its size in the context. The permutation expression is a non-empty vector of integers. The elements of the permutation vector must be continuous from 0 and cannot be repeated (i.e. [1,2,0] is a valid permutation. [2,0] or [1,1,2] is an invalid invalid permutation). The affine map is owned by the context.
Creates a symbol reference attribute in the given context referencing a symbol identified by the given string inside a list of nested references. Each of the references in the list must not be nested.
Creates a tensor type of a fixed rank with the given shape, element type, and optional encoding in the same context as the element type. The type is owned by the context. Tensor types without any specific encoding field should assign mlirAttributeGetNull to this parameter. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an unranked tensor type with the given element type in the same context as the element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates a vector type of the shape identified by its rank and dimensions, with the given element type in the same context as the element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Add a pass and transfer ownership to the provided OpPassManager. If the pass is not a generic operation pass or matching the type of the provided OpPassManager, a new OpPassManager is implicitly nested under the provided OpPassManager.
Add a pass and transfer ownership to the provided top-level PassManager. If the pass is not a generic operation pass or a ModulePass, a new OpPassManager is implicitly nested under the provided PassManager.
Parse a sequence of textual MLIR pass pipeline elements and add them to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Takes an operation owned by the caller and inserts it after the (non-owned) reference operation in the given block. If the reference is null, prepends the operation. Otherwise, the reference must belong to the block.
Takes a block owned by the caller and inserts it after the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, prepends the block to the region.
Takes an operation owned by the caller and inserts it before the (non-owned) reference operation in the given block. If the reference is null, appends the operation. Otherwise, the reference must belong to the block.
Takes a block owned by the caller and inserts it before the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, appends the block to the region.
Returns whether the given fully-qualified operation (i.e. 'dialect.operation') is registered with the context. This will return true if the dialect is loaded and the operation is registered within the dialect.
Checks whether the given affine map is an identity affine map. The function asserts that the number of dimensions is greater or equal to the number of results.
Looks up a symbol with the given name in the given symbol table and returns the operation that corresponds to the symbol. If the symbol cannot be found, returns a null operation.
Returns the affine map consisting of the most major nresults results. Returns the null AffineMap if the nresults is equal to zero. Returns the affineMap if nresults is greater or equals to number of results of the given affine map.
Returns the affine map consisting of the most minor nresults results. Returns the null AffineMap if the nresults is equal to zero. Returns the affineMap if nresults is greater or equals to number of results of the given affine map.
Moves the given operation immediately after the other operation in its parent block. The given operation may be owned by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Moves the given operation immediately before the other operation in its parent block. The given operation may be owner by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
`,6))]),t[3184]||(t[3184]=e("h1",{id:"MLIR-C-API",tabindex:"-1"},[s("MLIR C API "),e("a",{class:"header-anchor",href:"#MLIR-C-API","aria-label":'Permalink to "MLIR C API {#MLIR-C-API}"'},"")],-1)),e("details",Bi,[e("summary",null,[t[884]||(t[884]=e("a",{id:"Reactant.MLIR.API.LLVMAttributeRef",href:"#Reactant.MLIR.API.LLVMAttributeRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMAttributeRef")],-1)),t[885]||(t[885]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[886]||(t[886]=e("p",null,"Used to represent an attributes.",-1)),t[887]||(t[887]=e("p",null,[e("strong",null,"See also")],-1)),t[888]||(t[888]=e("p",null,"llvm::Attribute",-1)),t[889]||(t[889]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9761-L9766",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Gi,[e("summary",null,[t[890]||(t[890]=e("a",{id:"Reactant.MLIR.API.LLVMBasicBlockRef",href:"#Reactant.MLIR.API.LLVMBasicBlockRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBasicBlockRef")],-1)),t[891]||(t[891]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[892]||(t[892]=e("p",null,"Represents a basic block of instructions in LLVM IR.",-1)),t[893]||(t[893]=e("p",null,"This models llvm::BasicBlock.",-1)),t[894]||(t[894]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9674-L9678",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",zi,[e("summary",null,[t[895]||(t[895]=e("a",{id:"Reactant.MLIR.API.LLVMBinaryRef",href:"#Reactant.MLIR.API.LLVMBinaryRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBinaryRef")],-1)),t[896]||(t[896]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[897]||(t[897]=e("p",null,[e("strong",null,"See also")],-1)),t[898]||(t[898]=e("p",null,"llvm::object::Binary",-1)),t[899]||(t[899]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9803-L9806",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",wi,[e("summary",null,[t[900]||(t[900]=e("a",{id:"Reactant.MLIR.API.LLVMBool",href:"#Reactant.MLIR.API.LLVMBool"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBool")],-1)),t[901]||(t[901]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[902]||(t[902]=e("p",null,[e("code",null,"LLVMCSupportTypes Types and Enumerations")],-1)),t[903]||(t[903]=e("p",null,"@{",-1)),t[904]||(t[904]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9619-L9623",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Si,[e("summary",null,[t[905]||(t[905]=e("a",{id:"Reactant.MLIR.API.LLVMBuilderRef",href:"#Reactant.MLIR.API.LLVMBuilderRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBuilderRef")],-1)),t[906]||(t[906]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[907]||(t[907]=e("p",null,"Represents an LLVM basic block builder.",-1)),t[908]||(t[908]=e("p",null,"This models llvm::IRBuilder.",-1)),t[909]||(t[909]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9710-L9714",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ni,[e("summary",null,[t[910]||(t[910]=e("a",{id:"Reactant.MLIR.API.LLVMComdatRef",href:"#Reactant.MLIR.API.LLVMComdatRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMComdatRef")],-1)),t[911]||(t[911]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[912]||(t[912]=e("p",null,[e("strong",null,"See also")],-1)),t[913]||(t[913]=e("p",null,"llvm::Comdat",-1)),t[914]||(t[914]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9779-L9782",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Vi,[e("summary",null,[t[915]||(t[915]=e("a",{id:"Reactant.MLIR.API.LLVMContextRef",href:"#Reactant.MLIR.API.LLVMContextRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMContextRef")],-1)),t[916]||(t[916]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[917]||(t[917]=e("p",null,"The top-level container for all LLVM global data. See the LLVMContext class.",-1)),t[918]||(t[918]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9638-L9640",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",qi,[e("summary",null,[t[919]||(t[919]=e("a",{id:"Reactant.MLIR.API.LLVMDIBuilderRef",href:"#Reactant.MLIR.API.LLVMDIBuilderRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMDIBuilderRef")],-1)),t[920]||(t[920]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[921]||(t[921]=e("p",null,"Represents an LLVM debug info builder.",-1)),t[922]||(t[922]=e("p",null,"This models llvm::DIBuilder.",-1)),t[923]||(t[923]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9719-L9723",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ui,[e("summary",null,[t[924]||(t[924]=e("a",{id:"Reactant.MLIR.API.LLVMDbgRecordRef",href:"#Reactant.MLIR.API.LLVMDbgRecordRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMDbgRecordRef")],-1)),t[925]||(t[925]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[926]||(t[926]=e("p",null,[e("strong",null,"See also")],-1)),t[927]||(t[927]=e("p",null,"llvm::DbgRecord",-1)),t[928]||(t[928]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9811-L9814",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Qi,[e("summary",null,[t[929]||(t[929]=e("a",{id:"Reactant.MLIR.API.LLVMDiagnosticInfoRef",href:"#Reactant.MLIR.API.LLVMDiagnosticInfoRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMDiagnosticInfoRef")],-1)),t[930]||(t[930]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[931]||(t[931]=e("p",null,[e("strong",null,"See also")],-1)),t[932]||(t[932]=e("p",null,"llvm::DiagnosticInfo",-1)),t[933]||(t[933]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9771-L9774",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Wi,[e("summary",null,[t[934]||(t[934]=e("a",{id:"Reactant.MLIR.API.LLVMJITEventListenerRef",href:"#Reactant.MLIR.API.LLVMJITEventListenerRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMJITEventListenerRef")],-1)),t[935]||(t[935]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[936]||(t[936]=e("p",null,[e("strong",null,"See also")],-1)),t[937]||(t[937]=e("p",null,"llvm::JITEventListener",-1)),t[938]||(t[938]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9795-L9798",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Hi,[e("summary",null,[t[939]||(t[939]=e("a",{id:"Reactant.MLIR.API.LLVMMemoryBufferRef",href:"#Reactant.MLIR.API.LLVMMemoryBufferRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMMemoryBufferRef")],-1)),t[940]||(t[940]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[941]||(t[941]=e("p",null,"Used to pass regions of memory through LLVM interfaces.",-1)),t[942]||(t[942]=e("p",null,[e("strong",null,"See also")],-1)),t[943]||(t[943]=e("p",null,"llvm::MemoryBuffer",-1)),t[944]||(t[944]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9628-L9633",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Zi,[e("summary",null,[t[945]||(t[945]=e("a",{id:"Reactant.MLIR.API.LLVMMetadataRef",href:"#Reactant.MLIR.API.LLVMMetadataRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMMetadataRef")],-1)),t[946]||(t[946]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[947]||(t[947]=e("p",null,"Represents an LLVM Metadata.",-1)),t[948]||(t[948]=e("p",null,"This models llvm::Metadata.",-1)),t[949]||(t[949]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9683-L9687",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ji,[e("summary",null,[t[950]||(t[950]=e("a",{id:"Reactant.MLIR.API.LLVMModuleFlagEntry",href:"#Reactant.MLIR.API.LLVMModuleFlagEntry"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMModuleFlagEntry")],-1)),t[951]||(t[951]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[952]||(t[952]=e("p",null,[e("strong",null,"See also")],-1)),t[953]||(t[953]=e("p",null,"llvm::Module::ModuleFlagEntry",-1)),t[954]||(t[954]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9787-L9790",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ki,[e("summary",null,[t[955]||(t[955]=e("a",{id:"Reactant.MLIR.API.LLVMModuleProviderRef",href:"#Reactant.MLIR.API.LLVMModuleProviderRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMModuleProviderRef")],-1)),t[956]||(t[956]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[957]||(t[957]=e("p",null,"Interface used to provide a module to JIT or interpreter. This is now just a synonym for llvm::Module, but we have to keep using the different type to keep binary compatibility.",-1)),t[958]||(t[958]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9728-L9730",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",$i,[e("summary",null,[t[959]||(t[959]=e("a",{id:"Reactant.MLIR.API.LLVMModuleRef",href:"#Reactant.MLIR.API.LLVMModuleRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMModuleRef")],-1)),t[960]||(t[960]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[961]||(t[961]=e("p",null,"The top-level container for all other LLVM Intermediate Representation (IR) objects.",-1)),t[962]||(t[962]=e("p",null,[e("strong",null,"See also")],-1)),t[963]||(t[963]=e("p",null,"llvm::Module",-1)),t[964]||(t[964]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9645-L9650",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Xi,[e("summary",null,[t[965]||(t[965]=e("a",{id:"Reactant.MLIR.API.LLVMNamedMDNodeRef",href:"#Reactant.MLIR.API.LLVMNamedMDNodeRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMNamedMDNodeRef")],-1)),t[966]||(t[966]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[967]||(t[967]=e("p",null,"Represents an LLVM Named Metadata Node.",-1)),t[968]||(t[968]=e("p",null,"This models llvm::NamedMDNode.",-1)),t[969]||(t[969]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9692-L9696",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Yi,[e("summary",null,[t[970]||(t[970]=e("a",{id:"Reactant.MLIR.API.LLVMOperandBundleRef",href:"#Reactant.MLIR.API.LLVMOperandBundleRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMOperandBundleRef")],-1)),t[971]||(t[971]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[972]||(t[972]=e("p",null,[e("strong",null,"See also")],-1)),t[973]||(t[973]=e("p",null,"llvm::OperandBundleDef",-1)),t[974]||(t[974]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9753-L9756",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",_i,[e("summary",null,[t[975]||(t[975]=e("a",{id:"Reactant.MLIR.API.LLVMPassManagerRef",href:"#Reactant.MLIR.API.LLVMPassManagerRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMPassManagerRef")],-1)),t[976]||(t[976]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[977]||(t[977]=e("p",null,[e("strong",null,"See also")],-1)),t[978]||(t[978]=e("p",null,"llvm::PassManagerBase",-1)),t[979]||(t[979]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9735-L9738",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",tl,[e("summary",null,[t[980]||(t[980]=e("a",{id:"Reactant.MLIR.API.LLVMTypeRef",href:"#Reactant.MLIR.API.LLVMTypeRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMTypeRef")],-1)),t[981]||(t[981]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[982]||(t[982]=e("p",null,[s("Each value in the LLVM IR has a type, an "),e("a",{href:"/Reactant.jl/previews/PR363/api/mlirc#Reactant.MLIR.API.LLVMTypeRef"},[e("code",null,"LLVMTypeRef")]),s(".")],-1)),t[983]||(t[983]=e("p",null,[e("strong",null,"See also")],-1)),t[984]||(t[984]=e("p",null,"llvm::Type",-1)),t[985]||(t[985]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9655-L9660",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",el,[e("summary",null,[t[986]||(t[986]=e("a",{id:"Reactant.MLIR.API.LLVMUseRef",href:"#Reactant.MLIR.API.LLVMUseRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMUseRef")],-1)),t[987]||(t[987]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[988]||(t[988]=e("p",null,"Used to get the users and usees of a Value.",-1)),t[989]||(t[989]=e("p",null,[e("strong",null,"See also")],-1)),t[990]||(t[990]=e("p",null,"llvm::Use",-1)),t[991]||(t[991]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9743-L9748",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",sl,[e("summary",null,[t[992]||(t[992]=e("a",{id:"Reactant.MLIR.API.LLVMValueMetadataEntry",href:"#Reactant.MLIR.API.LLVMValueMetadataEntry"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMValueMetadataEntry")],-1)),t[993]||(t[993]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[994]||(t[994]=e("p",null,"Represents an entry in a Global Object's metadata attachments.",-1)),t[995]||(t[995]=e("p",null,"This models std::pair",-1)),t[996]||(t[996]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9701-L9705",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",al,[e("summary",null,[t[997]||(t[997]=e("a",{id:"Reactant.MLIR.API.LLVMValueRef",href:"#Reactant.MLIR.API.LLVMValueRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMValueRef")],-1)),t[998]||(t[998]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[999]||(t[999]=e("p",null,"Represents an individual value in LLVM IR.",-1)),t[1e3]||(t[1e3]=e("p",null,"This models llvm::Value.",-1)),t[1001]||(t[1001]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9665-L9669",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",il,[e("summary",null,[t[1002]||(t[1002]=e("a",{id:"Reactant.MLIR.API.MlirDiagnostic",href:"#Reactant.MLIR.API.MlirDiagnostic"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnostic")],-1)),t[1003]||(t[1003]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1004]||(t[1004]=l('
julia
MlirDiagnostic
An opaque reference to a diagnostic, always owned by the diagnostics engine (context). Must not be stored outside of the diagnostic handler.
',3))]),e("details",ll,[e("summary",null,[t[1005]||(t[1005]=e("a",{id:"Reactant.MLIR.API.MlirDiagnosticHandler",href:"#Reactant.MLIR.API.MlirDiagnosticHandler"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnosticHandler")],-1)),t[1006]||(t[1006]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1007]||(t[1007]=e("p",null,[s("Diagnostic handler type. Accepts a reference to a diagnostic, which is only guaranteed to be live during the call. The handler is passed the "),e("code",null,"userData"),s(" that was provided when the handler was attached to a context. If the handler processed the diagnostic completely, it is expected to return success. Otherwise, it is expected to return failure to indicate that other handlers should attempt to process the diagnostic.")],-1)),t[1008]||(t[1008]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L6718-L6720",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",nl,[e("summary",null,[t[1009]||(t[1009]=e("a",{id:"Reactant.MLIR.API.MlirDiagnosticHandlerID",href:"#Reactant.MLIR.API.MlirDiagnosticHandlerID"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnosticHandlerID")],-1)),t[1010]||(t[1010]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1011]||(t[1011]=e("p",null,"Opaque identifier of a diagnostic handler, useful to detach a handler.",-1)),t[1012]||(t[1012]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L6712-L6714",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",pl,[e("summary",null,[t[1013]||(t[1013]=e("a",{id:"Reactant.MLIR.API.MlirDiagnosticSeverity",href:"#Reactant.MLIR.API.MlirDiagnosticSeverity"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnosticSeverity")],-1)),t[1014]||(t[1014]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1015]||(t[1015]=l('
Structure of external MlirPass callbacks. All callbacks are required to be set unless otherwise specified.
Field
Note
construct
This callback is called from the pass is created. This is analogous to a C++ pass constructor.
destruct
This callback is called when the pass is destroyed This is analogous to a C++ pass destructor.
initialize
This callback is optional. The callback is called before the pass is run, allowing a chance to initialize any complex state necessary for running the pass. See Pass::initialize(MLIRContext *).
clone
This callback is called when the pass is cloned. See Pass::clonePass().
run
This callback is called when the pass is run. See Pass::runOnOperation().
A logical result value, essentially a boolean with named states. LLVM convention for using boolean values to designate success or failure of an operation is a moving target, so MLIR opted for an explicit class. Instances of MlirLogicalResult must only be inspected using the associated functions.
This class contains all the information necessary to construct the operation. It owns the MlirRegions it has pointers to and does not own anything else. By default, the state can be constructed from a name and location, the latter being also used to access the context, and has no other components. These components can be added progressively until the operation is constructed. Users are not expected to rely on the internals of this class and should use mlirOperationState* functions instead.
',4))]),e("details",ul,[e("summary",null,[t[1031]||(t[1031]=e("a",{id:"Reactant.MLIR.API.MlirOperationWalkCallback",href:"#Reactant.MLIR.API.MlirOperationWalkCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirOperationWalkCallback")],-1)),t[1032]||(t[1032]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1033]||(t[1033]=e("p",null,[s("Operation walker type. The handler is passed an (opaque) reference to an operation and a pointer to a "),e("code",null,"userData"),s(".")],-1)),t[1034]||(t[1034]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L1495-L1497",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",bl,[e("summary",null,[t[1035]||(t[1035]=e("a",{id:"Reactant.MLIR.API.MlirShapedTypeComponentsCallback",href:"#Reactant.MLIR.API.MlirShapedTypeComponentsCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirShapedTypeComponentsCallback")],-1)),t[1036]||(t[1036]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1037]||(t[1037]=e("p",null,"These callbacks are used to return multiple shaped type components from functions while transferring ownership to the caller. The first argument is the has rank boolean followed by the the rank and a pointer to the shape (if applicable). The next argument is the element type, then the attribute. The last argument is an opaque pointer forwarded to the callback by the caller. This callback will be called potentially multiple times for each shaped type components.",-1)),t[1038]||(t[1038]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9100-L9102",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",gl,[e("summary",null,[t[1039]||(t[1039]=e("a",{id:"Reactant.MLIR.API.MlirSparseTensorLevelType",href:"#Reactant.MLIR.API.MlirSparseTensorLevelType"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirSparseTensorLevelType")],-1)),t[1040]||(t[1040]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1041]||(t[1041]=e("p",null,"Dimension level types (and properties) that define sparse tensors. See the documentation in SparseTensorAttrDefs.td for their meaning.",-1)),t[1042]||(t[1042]=e("p",null,"These correspond to SparseTensorEncodingAttr::LevelType in the C++ API. If updating, keep them in sync and update the static_assert in the impl file.",-1)),t[1043]||(t[1043]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L8452-L8456",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",yl,[e("summary",null,[t[1044]||(t[1044]=e("a",{id:"Reactant.MLIR.API.MlirStringCallback",href:"#Reactant.MLIR.API.MlirStringCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirStringCallback")],-1)),t[1045]||(t[1045]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1046]||(t[1046]=e("p",null,"A callback for returning string references.",-1)),t[1047]||(t[1047]=e("p",null,[s("This function is called back by the functions that need to return a reference to the portion of the string with the following arguments: - an "),e("a",{href:"/Reactant.jl/previews/PR363/api/mlirc#Reactant.MLIR.API.MlirStringRef"},[e("code",null,"MlirStringRef")]),s(" representing the current portion of the string - a pointer to user data forwarded from the printing call.")],-1)),t[1048]||(t[1048]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L108-L112",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",ml,[e("summary",null,[t[1049]||(t[1049]=e("a",{id:"Reactant.MLIR.API.MlirStringRef",href:"#Reactant.MLIR.API.MlirStringRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirStringRef")],-1)),t[1050]||(t[1050]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1051]||(t[1051]=l('
julia
MlirStringRef
A pointer to a sized fragment of a string, not necessarily null-terminated. Does not own the underlying string. This is equivalent to llvm::StringRef.
',4))]),e("details",kl,[e("summary",null,[t[1052]||(t[1052]=e("a",{id:"Reactant.MLIR.API.MlirTypesCallback",href:"#Reactant.MLIR.API.MlirTypesCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirTypesCallback")],-1)),t[1053]||(t[1053]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1054]||(t[1054]=e("p",null,"These callbacks are used to return multiple types from functions while transferring ownership to the caller. The first argument is the number of consecutive elements pointed to by the second argument. The third argument is an opaque pointer forwarded to the callback by the caller.",-1)),t[1055]||(t[1055]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9052-L9054",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",fl,[e("summary",null,[t[1056]||(t[1056]=e("a",{id:"Reactant.MLIR.API.MlirWalkOrder",href:"#Reactant.MLIR.API.MlirWalkOrder"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirWalkOrder")],-1)),t[1057]||(t[1057]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1058]||(t[1058]=l('
This function parses the given arguments using the LLVM command line parser. Note that the only stable thing about this function is its signature; you cannot rely on any particular set of command line arguments being interpreted the same way across LLVM versions.
This function will search through all previously loaded dynamic libraries for the symbol symbolName. If it is found, the address of that symbol is returned. If not, null is returned.
Prints an affine expression by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Returns the simplified affine map resulting from dropping the symbols that do not appear in any of the individual maps in affineMaps. Asserts that all maps in affineMaps are normalized to the same number of dims and symbols. Takes a callback populateResult to fill the res container with value m at entry idx. This allows returning without worrying about ownership considerations.
Creates an affine map with results defined by the given list of affine expressions. The map resulting map also has the requested number of input dimensions and symbols, regardless of them being used in the results.
Returns the affine map consisting of the most major numResults results. Returns the null AffineMap if the numResults is equal to zero. Returns the affineMap if numResults is greater or equals to number of results of the given affine map.
Returns the affine map consisting of the most minor numResults results. Returns the null AffineMap if the numResults is equal to zero. Returns the affineMap if numResults is greater or equals to number of results of the given affine map.
Checks whether the given affine map is an identity affine map. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an identity affine map on the most minor dimensions in the context. The affine map is owned by the context. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an affine map with a permutation expression and its size in the context. The permutation expression is a non-empty vector of integers. The elements of the permutation vector must be continuous from 0 and cannot be repeated (i.e. [1,2,0] is a valid permutation. [2,0] or [1,1,2] is an invalid permutation.) The affine map is owned by the context.
Prints an affine map by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Apply AffineExpr::replace(map) to each of the results and return a new new AffineMap with the new results and the specified number of dims and symbols.
Creates an instance of AnyQuantizedType with the given parameters in the same context as storageType and returns it. The instance is owned by the context.
Prints an attribute by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Takes an operation owned by the caller and inserts it as pos to the block. This is an expensive operation that scans the block linearly, prefer insertBefore/After instead.
Takes an operation owned by the caller and inserts it after the (non-owned) reference operation in the given block. If the reference is null, prepends the operation. Otherwise, the reference must belong to the block.
Takes an operation owned by the caller and inserts it before the (non-owned) reference operation in the given block. If the reference is null, appends the operation. Otherwise, the reference must belong to the block.
Prints a block by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates an instance of CalibratedQuantizedType with the given parameters in the same context as expressedType and returns it. The instance is owned by the context.
Attaches the diagnostic handler to the context. Handlers are invoked in the reverse order of attachment until one of them processes the diagnostic completely. When a handler is invoked it is passed the userData that was provided when it was attached. If non-NULL, deleteUserData is called once the system no longer needs to call the handler (for instance after the handler is detached or the context is destroyed). Returns an identifier that can be used to detach the handler.
Gets the dialect instance owned by the given context using the dialect namespace to identify it, loads (i.e., constructs the instance of) the dialect if necessary. If the dialect is not registered with the context, returns null. Use mlirContextLoad<Name>Dialect to load an unregistered dialect.
Returns whether the given fully-qualified operation (i.e. 'dialect.operation') is registered with the context. This will return true if the dialect is loaded and the operation is registered within the dialect.
Sets the thread pool of the context explicitly, enabling multithreading in the process. This API should be used to avoid re-creating thread pools in long-running applications that perform multiple compilations, see the C++ documentation for MLIRContext for details.
Creates an external MlirPass that calls the supplied callbacks using the supplied userData. If opName is empty, the pass is a generic operation pass. Otherwise it is an operation pass specific to the specified pass name.
Creates a dense elements attribute with the given shaped type from elements of a specific type. Expects the element type of the shaped type to match the data element type.
Creates a dense elements attribute with the given Shaped type and elements populated from a packed, row-major opaque buffer of contents.
The format of the raw buffer is a densely packed array of values that can be bitcast to the storage format of the element type specified. Types that are not byte aligned will be: - For bitwidth > 1: Rounded up to the next byte. - For bitwidth = 1: Packed into 8bit bytes with bits corresponding to the linear order of the shape type from MSB to LSB, padded to on the right.
A raw buffer of a single element (or for 1-bit, a byte of value 0 or 255) will be interpreted as a splat. User code should be prepared for additional, conformant patterns to be identified as splats in the future.
Creates a dense elements attribute that has the same data as the given dense elements attribute and a different shaped type. The new type must have the same total number of elements.
Gets the total number of elements in the given elements attribute. In order to iterate over the attribute, obtain its type, which must be a statically shaped type and use its sizes to build a multi-dimensional index.
Creates an ExecutionEngine for the provided ModuleOp. The ModuleOp is expected to be "translatable" to LLVM IR (only contains operations in dialects that implement the LLVMTranslationDialectInterface). The module ownership stays with the client and can be destroyed as soon as the call returns. optLevel is the optimization level to be used for transformation and code generation. LLVM passes at optLevel are run before code generation. The number and array of paths corresponding to shared libraries that will be loaded are specified via numPaths and sharedLibPaths respectively. TODO: figure out other options.
Invoke a native function in the execution engine by name with the arguments and result of the invoked function passed as an array of pointers. The function must have been tagged with the llvm.emit\\_c\\_interface attribute. Returns a failure if the execution fails for any reason (the function name can't be resolved for instance).
Infers the return types of the operation identified by its canonical given the arguments that will be supplied to its generic builder. Calls callback with the types of inferred arguments, potentially several times, on success. Returns failure otherwise.
Checks if two integer set objects are equal. This is a "shallow" comparison of two objects. Only the sets with some small number of constraints are uniqued and compare equal here. Set objects that represent the same integer set with different constraints may be considered non-equal by this check. Set difference followed by an (expensive) emptiness check should be used to check equivalence of the underlying integer sets.
Gets or creates a new integer set in the given context. The set is defined by a list of affine constraints, with the given number of input dimensions and symbols, which are treated as either equalities (eqFlags is 1) or inequalities (eqFlags is 0). Both constraints and eqFlags are expected to point to at least numConstraint consecutive values.
Prints an integer set by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Gets or creates a new integer set in which the values and dimensions of the given set are replaced with the given affine expressions. dimReplacements and symbolReplacements are expected to point to at least as many consecutive expressions as the given set has dimensions and symbols, respectively. The new set will have numResultDims and numResultSymbols dimensions and symbols, respectively.
Creates a LLVM DIDerivedType attribute. Note that dwarfAddressSpace is an optional field, where MLIR_CAPI_DWARF_ADDRESS_SPACE_NULL indicates null and non-negative values indicate a value present.
Creates an LLVM identified struct type with no body. If a struct type with this name already exists in the context, returns that type. Use mlirLLVMStructTypeIdentifiedNewGet to create a fresh struct type, potentially renaming it. The body should be set separatelty by calling mlirLLVMStructTypeSetBody, if it isn't set already.
Creates an LLVM identified struct type with no body and a name starting with the given prefix. If a struct with the exact name as the given prefix already exists, appends an unspecified suffix to the name so that the name is unique in context.
Creates an LLVM literal (unnamed) struct type. This may assert if the fields have types not compatible with the LLVM dialect. For a graceful failure, use the checked version.
Creates a name location owned by the given context. Providing null location for childLoc is allowed and if childLoc is null location, then the behavior is the same as having unknown child location.
Prints a location by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates a MemRef type with the given rank, shape, memory space and element type in the same context as the element type. The type has no affine maps, i.e. represents a default row-major contiguous memref. The type is owned by the context.
Creates a MemRef type with the given rank and shape, a potentially empty list of affine layout maps, the given memory space and element type, in the same context as element type. The type is owned by the context.
Returns the strides of the MemRef if the layout map is in strided form. Both strides and offset are out params. strides must point to pre-allocated memory of length equal to the rank of the memref.
Merge the symbols from other into target, potentially renaming them to avoid conflicts. Private symbols may be renamed during the merge, public symbols must have at most one declaration. A name conflict in public symbols is reported as an error before returning a failure.
Note that this clones the other operation unlike the C++ counterpart that takes ownership.
Add a pass and transfer ownership to the provided mlirOpPassManager. If the pass is not a generic operation pass or matching the type of the provided PassManager, a new OpPassManager is implicitly nested under the provided PassManager.
Parse a sequence of textual MLIR pass pipeline elements and add them to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Nest an OpPassManager under the provided OpPassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed.
Enables the elision of large elements attributes by printing a lexically valid but otherwise meaningless form instead of the element data. The largeElementLimit is used to configure what is considered to be a "large" ElementsAttr by providing an upper limit to the number of elements.
Enables the elision of large resources strings by omitting them from the dialect_resources section. The largeResourceLimit is used to configure what is considered to be a "large" resource by providing an upper limit to the string size.
Enable or disable printing of debug information (based on enable). If 'prettyForm' is set to true, debug information is printed in a more readable 'pretty' form. Note: The IR generated with 'prettyForm' is not parsable.
Use local scope when printing the operation. This allows for using the printer in a more localized and thread-safe setting, but may not necessarily be identical to what the IR will look like when dumping the full module.
Creates an opaque attribute in the given context associated with the dialect identified by its namespace. The attribute contains opaque byte data of the specified length (data need not be null-terminated).
Creates an opaque type in the given context associated with the dialect identified by its namespace. The type contains opaque byte data of the specified length (data need not be null-terminated).
Creates an operation and transfers ownership to the caller. Note that caller owned child objects are transferred in this call and must not be further used. Particularly, this applies to any regions added to the state (the implementation may invalidate any such pointers).
This call can fail under the following conditions, in which case, it will return a null operation and emit diagnostics: - Result type inference is enabled and cannot be performed.
Parses an operation, giving ownership to the caller. If parsing fails a null operation will be returned, and an error diagnostic emitted.
sourceStr may be either the text assembly format, or binary bytecode format. sourceName is used as the file name of the source; any IR without locations will get a FileLineColLoc location with sourceName as the file name.
Returns the number of attributes attached to the operation. Deprecated, please use mlirOperationGetNumInherentAttributes or mlirOperationGetNumDiscardableAttributes.
Returns true if this operation defines an inherent attribute with this name. Note: the attribute can be optional, so mlirOperationGetInherentAttributeByName can still return a null attribute.
Returns true if the operation identified by its canonical string name implements the interface identified by its TypeID in the given context. Note that interfaces may be attached to operations in some contexts and not others.
Moves the given operation immediately after the other operation in its parent block. The given operation may be owned by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Moves the given operation immediately before the other operation in its parent block. The given operation may be owner by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Prints an operation by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Removes an attribute by name. Returns false if the attribute was not found and true if removed. Deprecated, please use mlirOperationRemoveInherentAttributeByName or mlirOperationRemoveDiscardableAttributeByName.
Sets a discardable attribute by name, replacing the existing if it exists or adding a new one otherwise. The new attr Attribute is not allowed to be null, use mlirOperationRemoveDiscardableAttributeByName to remove an Attribute instead.
Enables result type inference for the operation under construction. If enabled, then the caller must not have called mlirOperationStateAddResults(). Note that if enabled, the mlirOperationCreate() call is failable: it will return a null operation on inference failure and will emit diagnostics.
Walks operation op in walkOrder and calls callback on that operation. *userData is passed to the callback as well and can be used to tunnel some context or other data into the callback.
Parse a textual MLIR pass pipeline and assign it to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Add a pass and transfer ownership to the provided top-level mlirPassManager. If the pass is not a generic operation pass or a ModulePass, a new OpPassManager is implicitly nested under the provided PassManager.
Nest an OpPassManager under the top-level PassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed. To further nest more OpPassManager under the newly returned one, see mlirOpPassManagerNest below.
Print a textual MLIR pass pipeline by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Casts from a type based on the expressed type of the given type to a corresponding type based on the given type. Returns a null type if the cast is not valid.
Casts from a type based on the storage type of the given type to a corresponding type based on the given type. Returns a null type if the cast is not valid.
Creates a tensor type of a fixed rank with the given shape, element type, and optional encoding in the same context as the element type. The type is owned by the context. Tensor types without any specific encoding field should assign mlirAttributeGetNull() to this parameter.
Takes a block owned by the caller and inserts it at pos to the given region. This is an expensive operation that linearly scans the region, prefer insertAfter/Before instead.
Takes a block owned by the caller and inserts it after the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, prepends the block to the region.
Takes a block owned by the caller and inserts it before the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, appends the block to the region.
Reset the insertion point to no location. Creating an operation without a set insertion point is an error, but this can still be useful when the current insertion point a builder refers to is being removed.
Add new block with 'argTypes' arguments and set the insertion point to the end of it. The block is placed before 'insertBefore'. locs contains the locations of the inserted arguments, and should match the size of argTypes.
This method is used to signal the end of an in-place modification of the given operation. This can only be called on operations that were provided to a call to startOpModification.
Inline the operations of block 'source' before the operation 'op'. The source block will be deleted and must have no uses. 'argValues' is used to replace the block arguments of 'source'
The source block must have no successors. Otherwise, the resulting IR would have unreachable operations.
Move the blocks that belong to "region" before the given position in another region "parent". The two regions must be different. The caller is responsible for creating or updating the operation transferring flow of control to the region and passing it the correct block arguments.
Inline the operations of block 'source' into the end of block 'dest'. The source block will be deleted and must have no uses. 'argValues' is used to replace the block arguments of 'source'
The dest block must have no successors. Otherwise, the resulting IR would have unreachable operation.
Find uses of from and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced) and that the from operation is about to be replaced.
Find uses of from and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced) and that the from operation is about to be replaced.
mlirRewriterBaseReplaceAllUsesExcept(rewriter, from, to, exceptedUser)
Find uses of from and replace them with to except if the user is exceptedUser. Also notify the listener about every in-place op modification (for every use that was replaced).
Find uses of from within block and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced). The optional allUsesReplaced flag is set to "true" if all uses were replaced.
Replace the results of the given (original) operation with the specified new op (replacement). The result types of the two ops must match. The original op is erased.
Replace the results of the given (original) operation with the specified list of values (replacements). The result types of the given op and the replacements must match. The original op is erased.
Sets the insertion point to the node after the specified value. If value has a defining operation, sets the insertion point to the node after such defining operation. This will cause subsequent insertions to go right after it. Otherwise, value is a BlockArgument. Sets the insertion point to the start of its block.
This method is used to notify the rewriter that an in-place operation modification is about to happen. A call to this function must be followed by a call to either finalizeOpModification or cancelOpModification. This is a minor efficiency win (it avoids creating a new operation and removing the old one) but also often allows simpler code in the client.
Sets the current debug type, similarly to -debug-only=type in the command-line tools. Note that global debug should be enabled for any output to be produced.
Sets multiple current debug types, similarly to `-debug-only=type1,type2" in the command-line tools. Note that global debug should be enabled for any output to be produced.
Returns the value indicating a dynamic stride or offset in a shaped type. Prefer mlirShapedTypeGetDynamicStrideOrOffset to direct comparisons with this value.
Creates a sparse elements attribute of the given shape from a list of indices and a list of associated values. Both lists are expected to be dense elements attributes with the same number of elements. The list of indices is expected to contain 64-bit integers. The attribute is created in the same context as the type.
Creates a symbol reference attribute in the given context referencing a symbol identified by the given string inside a list of nested references. Each of the references in the list must not be nested.
Inserts the given operation into the given symbol table. The operation must have the symbol trait. If the symbol table already has a symbol with the same name, renames the symbol being inserted to ensure name uniqueness. Note that this does not move the operation itself into the block of the symbol table operation, this should be done separately. Returns the name of the symbol after insertion.
Looks up a symbol with the given name in the given symbol table and returns the operation that corresponds to the symbol. If the symbol cannot be found, returns a null operation.
Attempt to replace all uses that are nested within the given operation of the given symbol 'oldSymbol' with the provided 'newSymbol'. This does not traverse into nested symbol tables. Will fail atomically if there are any unknown operations that may be potential symbol tables.
Walks all symbol table operations nested within, and including, op. For each symbol table operation, the provided callback is invoked with the op and a boolean signifying if the symbols within that symbol table can be treated as if all uses within the IR are visible to the caller. allSymUsesVisible identifies whether all of the symbol uses of symbols within op are visible.
Applies the transformation script starting at the given transform root operation to the given payload operation. The module containing the transform root as well as the transform options should be provided. The transform operation must implement TransformOpInterface and the module must be a ModuleOp. Returns the status of the application.
Translate operation that satisfies LLVM dialect module requirements into an LLVM IR module living in the given context. This translates operations from any dilalect that has a registered implementation of LLVMTranslationDialectInterface.
Returns
the generated LLVM IR Module from the translated MLIR module, it is owned by the caller.
Prints a location by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates an instance of UniformQuantizedPerAxisType with the given parameters in the same context as storageType and returns it. scales and zeroPoints point to nDims number of elements. The instance is owned by the context.
Creates an instance of UniformQuantizedType with the given parameters in the same context as storageType and returns it. The instance is owned by the context.
Unlike the typed accessors below, constructs the attribute with a raw data buffer and no type/alignment checking. Use a more strongly typed accessor if possible. If dataIsMutable is false, then an immutable AsmResourceBlob will be created and that passed data contents will be treated as const. If the deleter is non NULL, then it will be called when the data buffer can no longer be accessed (passing userData to it).
Prints a value by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Replace all uses of 'of' value with the 'with' value, updating anything in the IR that uses 'of' to use the other value instead. When this returns there are zero uses of 'of'.
Creates a vector type of the shape identified by its rank and dimensions, with the given element type in the same context as the element type. The type is owned by the context.
Creates a scalable vector type with the shape identified by its rank and dimensions. A subset of dimensions may be marked as scalable via the corresponding flag list, which is expected to have as many entries as the rank of the vector. The vector is created in the same context as the element type.
Creates an affine map with results defined by the given list of affine expressions. The map resulting map also has the requested number of input dimensions and symbols, regardless of them being used in the results.
',3))]),e("details",v,[e("summary",null,[t[48]||(t[48]=e("a",{id:"Reactant.MLIR.IR.Attribute-Tuple{T} where T<:AbstractFloat",href:"#Reactant.MLIR.IR.Attribute-Tuple{T} where T<:AbstractFloat"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Attribute")],-1)),t[49]||(t[49]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[50]||(t[50]=l('
Creates a floating point attribute in the given context with the given double value and double-precision FP semantics. If check=true, emits appropriate diagnostics on illegal arguments.
',3))]),e("details",T,[e("summary",null,[t[51]||(t[51]=e("a",{id:"Reactant.MLIR.IR.Attribute-Tuple{T} where T<:Complex",href:"#Reactant.MLIR.IR.Attribute-Tuple{T} where T<:Complex"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Attribute")],-1)),t[52]||(t[52]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[53]||(t[53]=l('
',3))]),e("details",F,[e("summary",null,[t[60]||(t[60]=e("a",{id:"Reactant.MLIR.IR.Attribute-Union{Tuple{T}, Tuple{T, Any}} where T<:Integer",href:"#Reactant.MLIR.IR.Attribute-Union{Tuple{T}, Tuple{T, Any}} where T<:Integer"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Attribute")],-1)),t[61]||(t[61]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[62]||(t[62]=l('
julia
Attribute(int)
Creates an integer attribute of the given type with the given integer value.
Creates an ExecutionEngine for the provided ModuleOp. The ModuleOp is expected to be "translatable" to LLVM IR (only contains operations in dialects that implement the LLVMTranslationDialectInterface). The module ownership stays with the client and can be destroyed as soon as the call returns. optLevel is the optimization level to be used for transformation and code generation. LLVM passes at optLevel are run before code generation. The number and array of paths corresponding to shared libraries that will be loaded are specified via numPaths and sharedLibPaths respectively. TODO: figure out other options.
Gets or creates a new integer set in the given context. The set is defined by a list of affine constraints, with the given number of input dimensions and symbols, which are treated as either equalities (eqflags is 1) or inequalities (eqflags is 0). Both constraints and eqflags need to be arrays of the same length.
A logical result value, essentially a boolean with named states. LLVM convention for using boolean values to designate success or failure of an operation is a moving target, so MLIR opted for an explicit class. Instances of LogicalResult must only be inspected using the associated functions.
Nest an OpPassManager under the provided OpPassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed.
Nest an OpPassManager under the top-level PassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed. To further nest more OpPassManager under the newly returned one, see mlirOpPassManagerNest below.
`,3))]),e("details",rt,[e("summary",null,[t[153]||(t[153]=e("a",{id:"Reactant.MLIR.IR.Type-Union{Tuple{Type{Complex{T}}}, Tuple{T}} where T",href:"#Reactant.MLIR.IR.Type-Union{Tuple{Type{Complex{T}}}, Tuple{T}} where T"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.Type")],-1)),t[154]||(t[154]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[155]||(t[155]=l('
julia
Type(Complex{T}) where {T}
Creates a complex type with the given element type in the same context as the element type. The type is owned by the context.
Checks if two integer set objects are equal. This is a "shallow" comparison of two objects. Only the sets with some small number of constraints are uniqued and compare equal here. Set objects that represent the same integer set with different constraints may be considered non-equal by this check. Set difference followed by an (expensive) emptiness check should be used to check equivalence of the underlying integer sets.
Takes an operation owned by the caller and inserts it as index to the block. This is an expensive operation that scans the block linearly, prefer insertBefore/After instead.
Takes a block owned by the caller and inserts it at index to the given region. This is an expensive operation that linearly scans the region, prefer insertAfter/Before instead.
Inserts the given operation into the given symbol table. The operation must have the symbol trait. If the symbol table already has a symbol with the same name, renames the symbol being inserted to ensure name uniqueness. Note that this does not move the operation itself into the block of the symbol table operation, this should be done separately. Returns the name of the symbol after insertion.
Apply AffineExpr::replace(map) to each of the results and return a new new AffineMap with the new results and the specified number of dims and symbols.
Gets or creates a new integer set in which the values and dimensions of the given set are replaced with the given affine expressions. dimReplacements and symbolReplacements are expected to point to at least as many consecutive expressions as the given set has dimensions and symbols, respectively. The new set will have numResultDims and numResultSymbols dimensions and symbols, respectively.
Creates a dense elements attribute that has the same data as the given dense elements attribute and a different shaped type. The new type must have the same total number of elements.
',3))]),e("details",Jt,[e("summary",null,[t[273]||(t[273]=e("a",{id:"Reactant.MLIR.IR.BFloat16Type-Tuple{}",href:"#Reactant.MLIR.IR.BFloat16Type-Tuple{}"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.BFloat16Type")],-1)),t[274]||(t[274]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[275]||(t[275]=e("p",null,"BFloat16Type(; context=context())",-1)),t[276]||(t[276]=e("p",null,"Creates a bf16 type in the given context. The type is owned by the context.",-1)),t[277]||(t[277]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/IR/Type.jl#L157-L161",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Kt,[e("summary",null,[t[278]||(t[278]=e("a",{id:"Reactant.MLIR.IR.ConstantAffineMap-Tuple{Any}",href:"#Reactant.MLIR.IR.ConstantAffineMap-Tuple{Any}"},[e("span",{class:"jlbinding"},"Reactant.MLIR.IR.ConstantAffineMap")],-1)),t[279]||(t[279]=s()),i(a,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),t[280]||(t[280]=l('
julia
ConstantAffineMap(val; context=context())
Creates a single constant result affine map in the context. The affine map is owned by the context.
Creates a dense elements attribute with the given shaped type from elements of a specific type. Expects the element type of the shaped type to match the data element type.
Creates a MemRef type with the given rank and shape, a potentially empty list of affine layout maps, the given memory space and element type, in the same context as element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates a MemRef type with the given rank, shape, memory space and element type in the same context as the element type. The type has no affine maps, i.e. represents a default row-major contiguous memref. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an Unranked MemRef type with the given element type and in the given memory space. The type is owned by the context of element type. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an identity affine map on the most minor dimensions in the context. The affine map is owned by the context. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an opaque attribute in the given context associated with the dialect identified by its namespace. The attribute contains opaque byte data of the specified length (data need not be null-terminated).
Creates an opaque type in the given context associated with the dialect identified by its namespace. The type contains opaque byte data of the specified length (data need not be null-terminated).
Creates an affine map with a permutation expression and its size in the context. The permutation expression is a non-empty vector of integers. The elements of the permutation vector must be continuous from 0 and cannot be repeated (i.e. [1,2,0] is a valid permutation. [2,0] or [1,1,2] is an invalid invalid permutation). The affine map is owned by the context.
Creates a symbol reference attribute in the given context referencing a symbol identified by the given string inside a list of nested references. Each of the references in the list must not be nested.
Creates a tensor type of a fixed rank with the given shape, element type, and optional encoding in the same context as the element type. The type is owned by the context. Tensor types without any specific encoding field should assign mlirAttributeGetNull to this parameter. If check=true, emits appropriate diagnostics on illegal arguments.
Creates an unranked tensor type with the given element type in the same context as the element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Creates a vector type of the shape identified by its rank and dimensions, with the given element type in the same context as the element type. The type is owned by the context. If check=true, emits appropriate diagnostics on illegal arguments.
Add a pass and transfer ownership to the provided OpPassManager. If the pass is not a generic operation pass or matching the type of the provided OpPassManager, a new OpPassManager is implicitly nested under the provided OpPassManager.
Add a pass and transfer ownership to the provided top-level PassManager. If the pass is not a generic operation pass or a ModulePass, a new OpPassManager is implicitly nested under the provided PassManager.
Parse a sequence of textual MLIR pass pipeline elements and add them to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Takes an operation owned by the caller and inserts it after the (non-owned) reference operation in the given block. If the reference is null, prepends the operation. Otherwise, the reference must belong to the block.
Takes a block owned by the caller and inserts it after the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, prepends the block to the region.
Takes an operation owned by the caller and inserts it before the (non-owned) reference operation in the given block. If the reference is null, appends the operation. Otherwise, the reference must belong to the block.
Takes a block owned by the caller and inserts it before the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, appends the block to the region.
Returns whether the given fully-qualified operation (i.e. 'dialect.operation') is registered with the context. This will return true if the dialect is loaded and the operation is registered within the dialect.
Checks whether the given affine map is an identity affine map. The function asserts that the number of dimensions is greater or equal to the number of results.
Looks up a symbol with the given name in the given symbol table and returns the operation that corresponds to the symbol. If the symbol cannot be found, returns a null operation.
Returns the affine map consisting of the most major nresults results. Returns the null AffineMap if the nresults is equal to zero. Returns the affineMap if nresults is greater or equals to number of results of the given affine map.
Returns the affine map consisting of the most minor nresults results. Returns the null AffineMap if the nresults is equal to zero. Returns the affineMap if nresults is greater or equals to number of results of the given affine map.
Moves the given operation immediately after the other operation in its parent block. The given operation may be owned by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Moves the given operation immediately before the other operation in its parent block. The given operation may be owner by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
`,6))]),t[3184]||(t[3184]=e("h1",{id:"MLIR-C-API",tabindex:"-1"},[s("MLIR C API "),e("a",{class:"header-anchor",href:"#MLIR-C-API","aria-label":'Permalink to "MLIR C API {#MLIR-C-API}"'},"")],-1)),e("details",Bi,[e("summary",null,[t[884]||(t[884]=e("a",{id:"Reactant.MLIR.API.LLVMAttributeRef",href:"#Reactant.MLIR.API.LLVMAttributeRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMAttributeRef")],-1)),t[885]||(t[885]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[886]||(t[886]=e("p",null,"Used to represent an attributes.",-1)),t[887]||(t[887]=e("p",null,[e("strong",null,"See also")],-1)),t[888]||(t[888]=e("p",null,"llvm::Attribute",-1)),t[889]||(t[889]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9761-L9766",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Gi,[e("summary",null,[t[890]||(t[890]=e("a",{id:"Reactant.MLIR.API.LLVMBasicBlockRef",href:"#Reactant.MLIR.API.LLVMBasicBlockRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBasicBlockRef")],-1)),t[891]||(t[891]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[892]||(t[892]=e("p",null,"Represents a basic block of instructions in LLVM IR.",-1)),t[893]||(t[893]=e("p",null,"This models llvm::BasicBlock.",-1)),t[894]||(t[894]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9674-L9678",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",zi,[e("summary",null,[t[895]||(t[895]=e("a",{id:"Reactant.MLIR.API.LLVMBinaryRef",href:"#Reactant.MLIR.API.LLVMBinaryRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBinaryRef")],-1)),t[896]||(t[896]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[897]||(t[897]=e("p",null,[e("strong",null,"See also")],-1)),t[898]||(t[898]=e("p",null,"llvm::object::Binary",-1)),t[899]||(t[899]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9803-L9806",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",wi,[e("summary",null,[t[900]||(t[900]=e("a",{id:"Reactant.MLIR.API.LLVMBool",href:"#Reactant.MLIR.API.LLVMBool"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBool")],-1)),t[901]||(t[901]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[902]||(t[902]=e("p",null,[e("code",null,"LLVMCSupportTypes Types and Enumerations")],-1)),t[903]||(t[903]=e("p",null,"@{",-1)),t[904]||(t[904]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9619-L9623",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Si,[e("summary",null,[t[905]||(t[905]=e("a",{id:"Reactant.MLIR.API.LLVMBuilderRef",href:"#Reactant.MLIR.API.LLVMBuilderRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMBuilderRef")],-1)),t[906]||(t[906]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[907]||(t[907]=e("p",null,"Represents an LLVM basic block builder.",-1)),t[908]||(t[908]=e("p",null,"This models llvm::IRBuilder.",-1)),t[909]||(t[909]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9710-L9714",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ni,[e("summary",null,[t[910]||(t[910]=e("a",{id:"Reactant.MLIR.API.LLVMComdatRef",href:"#Reactant.MLIR.API.LLVMComdatRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMComdatRef")],-1)),t[911]||(t[911]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[912]||(t[912]=e("p",null,[e("strong",null,"See also")],-1)),t[913]||(t[913]=e("p",null,"llvm::Comdat",-1)),t[914]||(t[914]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9779-L9782",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Vi,[e("summary",null,[t[915]||(t[915]=e("a",{id:"Reactant.MLIR.API.LLVMContextRef",href:"#Reactant.MLIR.API.LLVMContextRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMContextRef")],-1)),t[916]||(t[916]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[917]||(t[917]=e("p",null,"The top-level container for all LLVM global data. See the LLVMContext class.",-1)),t[918]||(t[918]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9638-L9640",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",qi,[e("summary",null,[t[919]||(t[919]=e("a",{id:"Reactant.MLIR.API.LLVMDIBuilderRef",href:"#Reactant.MLIR.API.LLVMDIBuilderRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMDIBuilderRef")],-1)),t[920]||(t[920]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[921]||(t[921]=e("p",null,"Represents an LLVM debug info builder.",-1)),t[922]||(t[922]=e("p",null,"This models llvm::DIBuilder.",-1)),t[923]||(t[923]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9719-L9723",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ui,[e("summary",null,[t[924]||(t[924]=e("a",{id:"Reactant.MLIR.API.LLVMDbgRecordRef",href:"#Reactant.MLIR.API.LLVMDbgRecordRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMDbgRecordRef")],-1)),t[925]||(t[925]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[926]||(t[926]=e("p",null,[e("strong",null,"See also")],-1)),t[927]||(t[927]=e("p",null,"llvm::DbgRecord",-1)),t[928]||(t[928]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9811-L9814",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Qi,[e("summary",null,[t[929]||(t[929]=e("a",{id:"Reactant.MLIR.API.LLVMDiagnosticInfoRef",href:"#Reactant.MLIR.API.LLVMDiagnosticInfoRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMDiagnosticInfoRef")],-1)),t[930]||(t[930]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[931]||(t[931]=e("p",null,[e("strong",null,"See also")],-1)),t[932]||(t[932]=e("p",null,"llvm::DiagnosticInfo",-1)),t[933]||(t[933]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9771-L9774",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Wi,[e("summary",null,[t[934]||(t[934]=e("a",{id:"Reactant.MLIR.API.LLVMJITEventListenerRef",href:"#Reactant.MLIR.API.LLVMJITEventListenerRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMJITEventListenerRef")],-1)),t[935]||(t[935]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[936]||(t[936]=e("p",null,[e("strong",null,"See also")],-1)),t[937]||(t[937]=e("p",null,"llvm::JITEventListener",-1)),t[938]||(t[938]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9795-L9798",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Hi,[e("summary",null,[t[939]||(t[939]=e("a",{id:"Reactant.MLIR.API.LLVMMemoryBufferRef",href:"#Reactant.MLIR.API.LLVMMemoryBufferRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMMemoryBufferRef")],-1)),t[940]||(t[940]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[941]||(t[941]=e("p",null,"Used to pass regions of memory through LLVM interfaces.",-1)),t[942]||(t[942]=e("p",null,[e("strong",null,"See also")],-1)),t[943]||(t[943]=e("p",null,"llvm::MemoryBuffer",-1)),t[944]||(t[944]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9628-L9633",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Zi,[e("summary",null,[t[945]||(t[945]=e("a",{id:"Reactant.MLIR.API.LLVMMetadataRef",href:"#Reactant.MLIR.API.LLVMMetadataRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMMetadataRef")],-1)),t[946]||(t[946]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[947]||(t[947]=e("p",null,"Represents an LLVM Metadata.",-1)),t[948]||(t[948]=e("p",null,"This models llvm::Metadata.",-1)),t[949]||(t[949]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9683-L9687",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ji,[e("summary",null,[t[950]||(t[950]=e("a",{id:"Reactant.MLIR.API.LLVMModuleFlagEntry",href:"#Reactant.MLIR.API.LLVMModuleFlagEntry"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMModuleFlagEntry")],-1)),t[951]||(t[951]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[952]||(t[952]=e("p",null,[e("strong",null,"See also")],-1)),t[953]||(t[953]=e("p",null,"llvm::Module::ModuleFlagEntry",-1)),t[954]||(t[954]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9787-L9790",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Ki,[e("summary",null,[t[955]||(t[955]=e("a",{id:"Reactant.MLIR.API.LLVMModuleProviderRef",href:"#Reactant.MLIR.API.LLVMModuleProviderRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMModuleProviderRef")],-1)),t[956]||(t[956]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[957]||(t[957]=e("p",null,"Interface used to provide a module to JIT or interpreter. This is now just a synonym for llvm::Module, but we have to keep using the different type to keep binary compatibility.",-1)),t[958]||(t[958]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9728-L9730",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",$i,[e("summary",null,[t[959]||(t[959]=e("a",{id:"Reactant.MLIR.API.LLVMModuleRef",href:"#Reactant.MLIR.API.LLVMModuleRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMModuleRef")],-1)),t[960]||(t[960]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[961]||(t[961]=e("p",null,"The top-level container for all other LLVM Intermediate Representation (IR) objects.",-1)),t[962]||(t[962]=e("p",null,[e("strong",null,"See also")],-1)),t[963]||(t[963]=e("p",null,"llvm::Module",-1)),t[964]||(t[964]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9645-L9650",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Xi,[e("summary",null,[t[965]||(t[965]=e("a",{id:"Reactant.MLIR.API.LLVMNamedMDNodeRef",href:"#Reactant.MLIR.API.LLVMNamedMDNodeRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMNamedMDNodeRef")],-1)),t[966]||(t[966]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[967]||(t[967]=e("p",null,"Represents an LLVM Named Metadata Node.",-1)),t[968]||(t[968]=e("p",null,"This models llvm::NamedMDNode.",-1)),t[969]||(t[969]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9692-L9696",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",Yi,[e("summary",null,[t[970]||(t[970]=e("a",{id:"Reactant.MLIR.API.LLVMOperandBundleRef",href:"#Reactant.MLIR.API.LLVMOperandBundleRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMOperandBundleRef")],-1)),t[971]||(t[971]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[972]||(t[972]=e("p",null,[e("strong",null,"See also")],-1)),t[973]||(t[973]=e("p",null,"llvm::OperandBundleDef",-1)),t[974]||(t[974]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9753-L9756",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",_i,[e("summary",null,[t[975]||(t[975]=e("a",{id:"Reactant.MLIR.API.LLVMPassManagerRef",href:"#Reactant.MLIR.API.LLVMPassManagerRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMPassManagerRef")],-1)),t[976]||(t[976]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[977]||(t[977]=e("p",null,[e("strong",null,"See also")],-1)),t[978]||(t[978]=e("p",null,"llvm::PassManagerBase",-1)),t[979]||(t[979]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9735-L9738",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",tl,[e("summary",null,[t[980]||(t[980]=e("a",{id:"Reactant.MLIR.API.LLVMTypeRef",href:"#Reactant.MLIR.API.LLVMTypeRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMTypeRef")],-1)),t[981]||(t[981]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[982]||(t[982]=e("p",null,[s("Each value in the LLVM IR has a type, an "),e("a",{href:"/Reactant.jl/previews/PR363/api/mlirc#Reactant.MLIR.API.LLVMTypeRef"},[e("code",null,"LLVMTypeRef")]),s(".")],-1)),t[983]||(t[983]=e("p",null,[e("strong",null,"See also")],-1)),t[984]||(t[984]=e("p",null,"llvm::Type",-1)),t[985]||(t[985]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9655-L9660",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",el,[e("summary",null,[t[986]||(t[986]=e("a",{id:"Reactant.MLIR.API.LLVMUseRef",href:"#Reactant.MLIR.API.LLVMUseRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMUseRef")],-1)),t[987]||(t[987]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[988]||(t[988]=e("p",null,"Used to get the users and usees of a Value.",-1)),t[989]||(t[989]=e("p",null,[e("strong",null,"See also")],-1)),t[990]||(t[990]=e("p",null,"llvm::Use",-1)),t[991]||(t[991]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9743-L9748",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",sl,[e("summary",null,[t[992]||(t[992]=e("a",{id:"Reactant.MLIR.API.LLVMValueMetadataEntry",href:"#Reactant.MLIR.API.LLVMValueMetadataEntry"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMValueMetadataEntry")],-1)),t[993]||(t[993]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[994]||(t[994]=e("p",null,"Represents an entry in a Global Object's metadata attachments.",-1)),t[995]||(t[995]=e("p",null,"This models std::pair",-1)),t[996]||(t[996]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9701-L9705",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",al,[e("summary",null,[t[997]||(t[997]=e("a",{id:"Reactant.MLIR.API.LLVMValueRef",href:"#Reactant.MLIR.API.LLVMValueRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.LLVMValueRef")],-1)),t[998]||(t[998]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[999]||(t[999]=e("p",null,"Represents an individual value in LLVM IR.",-1)),t[1e3]||(t[1e3]=e("p",null,"This models llvm::Value.",-1)),t[1001]||(t[1001]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9665-L9669",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",il,[e("summary",null,[t[1002]||(t[1002]=e("a",{id:"Reactant.MLIR.API.MlirDiagnostic",href:"#Reactant.MLIR.API.MlirDiagnostic"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnostic")],-1)),t[1003]||(t[1003]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1004]||(t[1004]=l('
julia
MlirDiagnostic
An opaque reference to a diagnostic, always owned by the diagnostics engine (context). Must not be stored outside of the diagnostic handler.
',3))]),e("details",ll,[e("summary",null,[t[1005]||(t[1005]=e("a",{id:"Reactant.MLIR.API.MlirDiagnosticHandler",href:"#Reactant.MLIR.API.MlirDiagnosticHandler"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnosticHandler")],-1)),t[1006]||(t[1006]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1007]||(t[1007]=e("p",null,[s("Diagnostic handler type. Accepts a reference to a diagnostic, which is only guaranteed to be live during the call. The handler is passed the "),e("code",null,"userData"),s(" that was provided when the handler was attached to a context. If the handler processed the diagnostic completely, it is expected to return success. Otherwise, it is expected to return failure to indicate that other handlers should attempt to process the diagnostic.")],-1)),t[1008]||(t[1008]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L6718-L6720",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",nl,[e("summary",null,[t[1009]||(t[1009]=e("a",{id:"Reactant.MLIR.API.MlirDiagnosticHandlerID",href:"#Reactant.MLIR.API.MlirDiagnosticHandlerID"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnosticHandlerID")],-1)),t[1010]||(t[1010]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1011]||(t[1011]=e("p",null,"Opaque identifier of a diagnostic handler, useful to detach a handler.",-1)),t[1012]||(t[1012]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L6712-L6714",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",pl,[e("summary",null,[t[1013]||(t[1013]=e("a",{id:"Reactant.MLIR.API.MlirDiagnosticSeverity",href:"#Reactant.MLIR.API.MlirDiagnosticSeverity"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirDiagnosticSeverity")],-1)),t[1014]||(t[1014]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1015]||(t[1015]=l('
Structure of external MlirPass callbacks. All callbacks are required to be set unless otherwise specified.
Field
Note
construct
This callback is called from the pass is created. This is analogous to a C++ pass constructor.
destruct
This callback is called when the pass is destroyed This is analogous to a C++ pass destructor.
initialize
This callback is optional. The callback is called before the pass is run, allowing a chance to initialize any complex state necessary for running the pass. See Pass::initialize(MLIRContext *).
clone
This callback is called when the pass is cloned. See Pass::clonePass().
run
This callback is called when the pass is run. See Pass::runOnOperation().
A logical result value, essentially a boolean with named states. LLVM convention for using boolean values to designate success or failure of an operation is a moving target, so MLIR opted for an explicit class. Instances of MlirLogicalResult must only be inspected using the associated functions.
This class contains all the information necessary to construct the operation. It owns the MlirRegions it has pointers to and does not own anything else. By default, the state can be constructed from a name and location, the latter being also used to access the context, and has no other components. These components can be added progressively until the operation is constructed. Users are not expected to rely on the internals of this class and should use mlirOperationState* functions instead.
',4))]),e("details",ul,[e("summary",null,[t[1031]||(t[1031]=e("a",{id:"Reactant.MLIR.API.MlirOperationWalkCallback",href:"#Reactant.MLIR.API.MlirOperationWalkCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirOperationWalkCallback")],-1)),t[1032]||(t[1032]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1033]||(t[1033]=e("p",null,[s("Operation walker type. The handler is passed an (opaque) reference to an operation and a pointer to a "),e("code",null,"userData"),s(".")],-1)),t[1034]||(t[1034]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L1495-L1497",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",bl,[e("summary",null,[t[1035]||(t[1035]=e("a",{id:"Reactant.MLIR.API.MlirShapedTypeComponentsCallback",href:"#Reactant.MLIR.API.MlirShapedTypeComponentsCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirShapedTypeComponentsCallback")],-1)),t[1036]||(t[1036]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1037]||(t[1037]=e("p",null,"These callbacks are used to return multiple shaped type components from functions while transferring ownership to the caller. The first argument is the has rank boolean followed by the the rank and a pointer to the shape (if applicable). The next argument is the element type, then the attribute. The last argument is an opaque pointer forwarded to the callback by the caller. This callback will be called potentially multiple times for each shaped type components.",-1)),t[1038]||(t[1038]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9100-L9102",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",gl,[e("summary",null,[t[1039]||(t[1039]=e("a",{id:"Reactant.MLIR.API.MlirSparseTensorLevelType",href:"#Reactant.MLIR.API.MlirSparseTensorLevelType"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirSparseTensorLevelType")],-1)),t[1040]||(t[1040]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1041]||(t[1041]=e("p",null,"Dimension level types (and properties) that define sparse tensors. See the documentation in SparseTensorAttrDefs.td for their meaning.",-1)),t[1042]||(t[1042]=e("p",null,"These correspond to SparseTensorEncodingAttr::LevelType in the C++ API. If updating, keep them in sync and update the static_assert in the impl file.",-1)),t[1043]||(t[1043]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L8452-L8456",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",yl,[e("summary",null,[t[1044]||(t[1044]=e("a",{id:"Reactant.MLIR.API.MlirStringCallback",href:"#Reactant.MLIR.API.MlirStringCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirStringCallback")],-1)),t[1045]||(t[1045]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1046]||(t[1046]=e("p",null,"A callback for returning string references.",-1)),t[1047]||(t[1047]=e("p",null,[s("This function is called back by the functions that need to return a reference to the portion of the string with the following arguments: - an "),e("a",{href:"/Reactant.jl/previews/PR363/api/mlirc#Reactant.MLIR.API.MlirStringRef"},[e("code",null,"MlirStringRef")]),s(" representing the current portion of the string - a pointer to user data forwarded from the printing call.")],-1)),t[1048]||(t[1048]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L108-L112",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",ml,[e("summary",null,[t[1049]||(t[1049]=e("a",{id:"Reactant.MLIR.API.MlirStringRef",href:"#Reactant.MLIR.API.MlirStringRef"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirStringRef")],-1)),t[1050]||(t[1050]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1051]||(t[1051]=l('
julia
MlirStringRef
A pointer to a sized fragment of a string, not necessarily null-terminated. Does not own the underlying string. This is equivalent to llvm::StringRef.
',4))]),e("details",kl,[e("summary",null,[t[1052]||(t[1052]=e("a",{id:"Reactant.MLIR.API.MlirTypesCallback",href:"#Reactant.MLIR.API.MlirTypesCallback"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirTypesCallback")],-1)),t[1053]||(t[1053]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1054]||(t[1054]=e("p",null,"These callbacks are used to return multiple types from functions while transferring ownership to the caller. The first argument is the number of consecutive elements pointed to by the second argument. The third argument is an opaque pointer forwarded to the callback by the caller.",-1)),t[1055]||(t[1055]=e("p",null,[e("a",{href:"https://github.com/EnzymeAD/Reactant.jl/blob/716daf609b6a2f5cd9031dd20a875c87fcf99c58/src/mlir/libMLIR_h.jl#L9052-L9054",target:"_blank",rel:"noreferrer"},"source")],-1))]),e("details",fl,[e("summary",null,[t[1056]||(t[1056]=e("a",{id:"Reactant.MLIR.API.MlirWalkOrder",href:"#Reactant.MLIR.API.MlirWalkOrder"},[e("span",{class:"jlbinding"},"Reactant.MLIR.API.MlirWalkOrder")],-1)),t[1057]||(t[1057]=s()),i(a,{type:"info",class:"jlObjectType jlType",text:"Type"})]),t[1058]||(t[1058]=l('
This function parses the given arguments using the LLVM command line parser. Note that the only stable thing about this function is its signature; you cannot rely on any particular set of command line arguments being interpreted the same way across LLVM versions.
This function will search through all previously loaded dynamic libraries for the symbol symbolName. If it is found, the address of that symbol is returned. If not, null is returned.
Prints an affine expression by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Returns the simplified affine map resulting from dropping the symbols that do not appear in any of the individual maps in affineMaps. Asserts that all maps in affineMaps are normalized to the same number of dims and symbols. Takes a callback populateResult to fill the res container with value m at entry idx. This allows returning without worrying about ownership considerations.
Creates an affine map with results defined by the given list of affine expressions. The map resulting map also has the requested number of input dimensions and symbols, regardless of them being used in the results.
Returns the affine map consisting of the most major numResults results. Returns the null AffineMap if the numResults is equal to zero. Returns the affineMap if numResults is greater or equals to number of results of the given affine map.
Returns the affine map consisting of the most minor numResults results. Returns the null AffineMap if the numResults is equal to zero. Returns the affineMap if numResults is greater or equals to number of results of the given affine map.
Checks whether the given affine map is an identity affine map. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an identity affine map on the most minor dimensions in the context. The affine map is owned by the context. The function asserts that the number of dimensions is greater or equal to the number of results.
Creates an affine map with a permutation expression and its size in the context. The permutation expression is a non-empty vector of integers. The elements of the permutation vector must be continuous from 0 and cannot be repeated (i.e. [1,2,0] is a valid permutation. [2,0] or [1,1,2] is an invalid permutation.) The affine map is owned by the context.
Prints an affine map by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Apply AffineExpr::replace(map) to each of the results and return a new new AffineMap with the new results and the specified number of dims and symbols.
Creates an instance of AnyQuantizedType with the given parameters in the same context as storageType and returns it. The instance is owned by the context.
Prints an attribute by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Takes an operation owned by the caller and inserts it as pos to the block. This is an expensive operation that scans the block linearly, prefer insertBefore/After instead.
Takes an operation owned by the caller and inserts it after the (non-owned) reference operation in the given block. If the reference is null, prepends the operation. Otherwise, the reference must belong to the block.
Takes an operation owned by the caller and inserts it before the (non-owned) reference operation in the given block. If the reference is null, appends the operation. Otherwise, the reference must belong to the block.
Prints a block by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates an instance of CalibratedQuantizedType with the given parameters in the same context as expressedType and returns it. The instance is owned by the context.
Attaches the diagnostic handler to the context. Handlers are invoked in the reverse order of attachment until one of them processes the diagnostic completely. When a handler is invoked it is passed the userData that was provided when it was attached. If non-NULL, deleteUserData is called once the system no longer needs to call the handler (for instance after the handler is detached or the context is destroyed). Returns an identifier that can be used to detach the handler.
Gets the dialect instance owned by the given context using the dialect namespace to identify it, loads (i.e., constructs the instance of) the dialect if necessary. If the dialect is not registered with the context, returns null. Use mlirContextLoad<Name>Dialect to load an unregistered dialect.
Returns whether the given fully-qualified operation (i.e. 'dialect.operation') is registered with the context. This will return true if the dialect is loaded and the operation is registered within the dialect.
Sets the thread pool of the context explicitly, enabling multithreading in the process. This API should be used to avoid re-creating thread pools in long-running applications that perform multiple compilations, see the C++ documentation for MLIRContext for details.
Creates an external MlirPass that calls the supplied callbacks using the supplied userData. If opName is empty, the pass is a generic operation pass. Otherwise it is an operation pass specific to the specified pass name.
Creates a dense elements attribute with the given shaped type from elements of a specific type. Expects the element type of the shaped type to match the data element type.
Creates a dense elements attribute with the given Shaped type and elements populated from a packed, row-major opaque buffer of contents.
The format of the raw buffer is a densely packed array of values that can be bitcast to the storage format of the element type specified. Types that are not byte aligned will be: - For bitwidth > 1: Rounded up to the next byte. - For bitwidth = 1: Packed into 8bit bytes with bits corresponding to the linear order of the shape type from MSB to LSB, padded to on the right.
A raw buffer of a single element (or for 1-bit, a byte of value 0 or 255) will be interpreted as a splat. User code should be prepared for additional, conformant patterns to be identified as splats in the future.
Creates a dense elements attribute that has the same data as the given dense elements attribute and a different shaped type. The new type must have the same total number of elements.
Gets the total number of elements in the given elements attribute. In order to iterate over the attribute, obtain its type, which must be a statically shaped type and use its sizes to build a multi-dimensional index.
Creates an ExecutionEngine for the provided ModuleOp. The ModuleOp is expected to be "translatable" to LLVM IR (only contains operations in dialects that implement the LLVMTranslationDialectInterface). The module ownership stays with the client and can be destroyed as soon as the call returns. optLevel is the optimization level to be used for transformation and code generation. LLVM passes at optLevel are run before code generation. The number and array of paths corresponding to shared libraries that will be loaded are specified via numPaths and sharedLibPaths respectively. TODO: figure out other options.
Invoke a native function in the execution engine by name with the arguments and result of the invoked function passed as an array of pointers. The function must have been tagged with the llvm.emit\\_c\\_interface attribute. Returns a failure if the execution fails for any reason (the function name can't be resolved for instance).
Infers the return types of the operation identified by its canonical given the arguments that will be supplied to its generic builder. Calls callback with the types of inferred arguments, potentially several times, on success. Returns failure otherwise.
Checks if two integer set objects are equal. This is a "shallow" comparison of two objects. Only the sets with some small number of constraints are uniqued and compare equal here. Set objects that represent the same integer set with different constraints may be considered non-equal by this check. Set difference followed by an (expensive) emptiness check should be used to check equivalence of the underlying integer sets.
Gets or creates a new integer set in the given context. The set is defined by a list of affine constraints, with the given number of input dimensions and symbols, which are treated as either equalities (eqFlags is 1) or inequalities (eqFlags is 0). Both constraints and eqFlags are expected to point to at least numConstraint consecutive values.
Prints an integer set by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Gets or creates a new integer set in which the values and dimensions of the given set are replaced with the given affine expressions. dimReplacements and symbolReplacements are expected to point to at least as many consecutive expressions as the given set has dimensions and symbols, respectively. The new set will have numResultDims and numResultSymbols dimensions and symbols, respectively.
Creates a LLVM DIDerivedType attribute. Note that dwarfAddressSpace is an optional field, where MLIR_CAPI_DWARF_ADDRESS_SPACE_NULL indicates null and non-negative values indicate a value present.
Creates an LLVM identified struct type with no body. If a struct type with this name already exists in the context, returns that type. Use mlirLLVMStructTypeIdentifiedNewGet to create a fresh struct type, potentially renaming it. The body should be set separatelty by calling mlirLLVMStructTypeSetBody, if it isn't set already.
Creates an LLVM identified struct type with no body and a name starting with the given prefix. If a struct with the exact name as the given prefix already exists, appends an unspecified suffix to the name so that the name is unique in context.
Creates an LLVM literal (unnamed) struct type. This may assert if the fields have types not compatible with the LLVM dialect. For a graceful failure, use the checked version.
Creates a name location owned by the given context. Providing null location for childLoc is allowed and if childLoc is null location, then the behavior is the same as having unknown child location.
Prints a location by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates a MemRef type with the given rank, shape, memory space and element type in the same context as the element type. The type has no affine maps, i.e. represents a default row-major contiguous memref. The type is owned by the context.
Creates a MemRef type with the given rank and shape, a potentially empty list of affine layout maps, the given memory space and element type, in the same context as element type. The type is owned by the context.
Returns the strides of the MemRef if the layout map is in strided form. Both strides and offset are out params. strides must point to pre-allocated memory of length equal to the rank of the memref.
Merge the symbols from other into target, potentially renaming them to avoid conflicts. Private symbols may be renamed during the merge, public symbols must have at most one declaration. A name conflict in public symbols is reported as an error before returning a failure.
Note that this clones the other operation unlike the C++ counterpart that takes ownership.
Add a pass and transfer ownership to the provided mlirOpPassManager. If the pass is not a generic operation pass or matching the type of the provided PassManager, a new OpPassManager is implicitly nested under the provided PassManager.
Parse a sequence of textual MLIR pass pipeline elements and add them to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Nest an OpPassManager under the provided OpPassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed.
Enables the elision of large elements attributes by printing a lexically valid but otherwise meaningless form instead of the element data. The largeElementLimit is used to configure what is considered to be a "large" ElementsAttr by providing an upper limit to the number of elements.
Enables the elision of large resources strings by omitting them from the dialect_resources section. The largeResourceLimit is used to configure what is considered to be a "large" resource by providing an upper limit to the string size.
Enable or disable printing of debug information (based on enable). If 'prettyForm' is set to true, debug information is printed in a more readable 'pretty' form. Note: The IR generated with 'prettyForm' is not parsable.
Use local scope when printing the operation. This allows for using the printer in a more localized and thread-safe setting, but may not necessarily be identical to what the IR will look like when dumping the full module.
Creates an opaque attribute in the given context associated with the dialect identified by its namespace. The attribute contains opaque byte data of the specified length (data need not be null-terminated).
Creates an opaque type in the given context associated with the dialect identified by its namespace. The type contains opaque byte data of the specified length (data need not be null-terminated).
Creates an operation and transfers ownership to the caller. Note that caller owned child objects are transferred in this call and must not be further used. Particularly, this applies to any regions added to the state (the implementation may invalidate any such pointers).
This call can fail under the following conditions, in which case, it will return a null operation and emit diagnostics: - Result type inference is enabled and cannot be performed.
Parses an operation, giving ownership to the caller. If parsing fails a null operation will be returned, and an error diagnostic emitted.
sourceStr may be either the text assembly format, or binary bytecode format. sourceName is used as the file name of the source; any IR without locations will get a FileLineColLoc location with sourceName as the file name.
Returns the number of attributes attached to the operation. Deprecated, please use mlirOperationGetNumInherentAttributes or mlirOperationGetNumDiscardableAttributes.
Returns true if this operation defines an inherent attribute with this name. Note: the attribute can be optional, so mlirOperationGetInherentAttributeByName can still return a null attribute.
Returns true if the operation identified by its canonical string name implements the interface identified by its TypeID in the given context. Note that interfaces may be attached to operations in some contexts and not others.
Moves the given operation immediately after the other operation in its parent block. The given operation may be owned by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Moves the given operation immediately before the other operation in its parent block. The given operation may be owner by the caller or by its current block. The other operation must belong to a block. In any case, the ownership is transferred to the block of the other operation.
Prints an operation by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Removes an attribute by name. Returns false if the attribute was not found and true if removed. Deprecated, please use mlirOperationRemoveInherentAttributeByName or mlirOperationRemoveDiscardableAttributeByName.
Sets a discardable attribute by name, replacing the existing if it exists or adding a new one otherwise. The new attr Attribute is not allowed to be null, use mlirOperationRemoveDiscardableAttributeByName to remove an Attribute instead.
Enables result type inference for the operation under construction. If enabled, then the caller must not have called mlirOperationStateAddResults(). Note that if enabled, the mlirOperationCreate() call is failable: it will return a null operation on inference failure and will emit diagnostics.
Walks operation op in walkOrder and calls callback on that operation. *userData is passed to the callback as well and can be used to tunnel some context or other data into the callback.
Parse a textual MLIR pass pipeline and assign it to the provided OpPassManager. If parsing fails an error message is reported using the provided callback.
Add a pass and transfer ownership to the provided top-level mlirPassManager. If the pass is not a generic operation pass or a ModulePass, a new OpPassManager is implicitly nested under the provided PassManager.
Nest an OpPassManager under the top-level PassManager, the nested passmanager will only run on operations matching the provided name. The returned OpPassManager will be destroyed when the parent is destroyed. To further nest more OpPassManager under the newly returned one, see mlirOpPassManagerNest below.
Print a textual MLIR pass pipeline by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Casts from a type based on the expressed type of the given type to a corresponding type based on the given type. Returns a null type if the cast is not valid.
Casts from a type based on the storage type of the given type to a corresponding type based on the given type. Returns a null type if the cast is not valid.
Creates a tensor type of a fixed rank with the given shape, element type, and optional encoding in the same context as the element type. The type is owned by the context. Tensor types without any specific encoding field should assign mlirAttributeGetNull() to this parameter.
Takes a block owned by the caller and inserts it at pos to the given region. This is an expensive operation that linearly scans the region, prefer insertAfter/Before instead.
Takes a block owned by the caller and inserts it after the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, prepends the block to the region.
Takes a block owned by the caller and inserts it before the (non-owned) reference block in the given region. The reference block must belong to the region. If the reference block is null, appends the block to the region.
Reset the insertion point to no location. Creating an operation without a set insertion point is an error, but this can still be useful when the current insertion point a builder refers to is being removed.
Add new block with 'argTypes' arguments and set the insertion point to the end of it. The block is placed before 'insertBefore'. locs contains the locations of the inserted arguments, and should match the size of argTypes.
This method is used to signal the end of an in-place modification of the given operation. This can only be called on operations that were provided to a call to startOpModification.
Inline the operations of block 'source' before the operation 'op'. The source block will be deleted and must have no uses. 'argValues' is used to replace the block arguments of 'source'
The source block must have no successors. Otherwise, the resulting IR would have unreachable operations.
Move the blocks that belong to "region" before the given position in another region "parent". The two regions must be different. The caller is responsible for creating or updating the operation transferring flow of control to the region and passing it the correct block arguments.
Inline the operations of block 'source' into the end of block 'dest'. The source block will be deleted and must have no uses. 'argValues' is used to replace the block arguments of 'source'
The dest block must have no successors. Otherwise, the resulting IR would have unreachable operation.
Find uses of from and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced) and that the from operation is about to be replaced.
Find uses of from and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced) and that the from operation is about to be replaced.
mlirRewriterBaseReplaceAllUsesExcept(rewriter, from, to, exceptedUser)
Find uses of from and replace them with to except if the user is exceptedUser. Also notify the listener about every in-place op modification (for every use that was replaced).
Find uses of from within block and replace them with to. Also notify the listener about every in-place op modification (for every use that was replaced). The optional allUsesReplaced flag is set to "true" if all uses were replaced.
Replace the results of the given (original) operation with the specified new op (replacement). The result types of the two ops must match. The original op is erased.
Replace the results of the given (original) operation with the specified list of values (replacements). The result types of the given op and the replacements must match. The original op is erased.
Sets the insertion point to the node after the specified value. If value has a defining operation, sets the insertion point to the node after such defining operation. This will cause subsequent insertions to go right after it. Otherwise, value is a BlockArgument. Sets the insertion point to the start of its block.
This method is used to notify the rewriter that an in-place operation modification is about to happen. A call to this function must be followed by a call to either finalizeOpModification or cancelOpModification. This is a minor efficiency win (it avoids creating a new operation and removing the old one) but also often allows simpler code in the client.
Sets the current debug type, similarly to -debug-only=type in the command-line tools. Note that global debug should be enabled for any output to be produced.
Sets multiple current debug types, similarly to `-debug-only=type1,type2" in the command-line tools. Note that global debug should be enabled for any output to be produced.
Returns the value indicating a dynamic stride or offset in a shaped type. Prefer mlirShapedTypeGetDynamicStrideOrOffset to direct comparisons with this value.
Creates a sparse elements attribute of the given shape from a list of indices and a list of associated values. Both lists are expected to be dense elements attributes with the same number of elements. The list of indices is expected to contain 64-bit integers. The attribute is created in the same context as the type.
Creates a symbol reference attribute in the given context referencing a symbol identified by the given string inside a list of nested references. Each of the references in the list must not be nested.
Inserts the given operation into the given symbol table. The operation must have the symbol trait. If the symbol table already has a symbol with the same name, renames the symbol being inserted to ensure name uniqueness. Note that this does not move the operation itself into the block of the symbol table operation, this should be done separately. Returns the name of the symbol after insertion.
Looks up a symbol with the given name in the given symbol table and returns the operation that corresponds to the symbol. If the symbol cannot be found, returns a null operation.
Attempt to replace all uses that are nested within the given operation of the given symbol 'oldSymbol' with the provided 'newSymbol'. This does not traverse into nested symbol tables. Will fail atomically if there are any unknown operations that may be potential symbol tables.
Walks all symbol table operations nested within, and including, op. For each symbol table operation, the provided callback is invoked with the op and a boolean signifying if the symbols within that symbol table can be treated as if all uses within the IR are visible to the caller. allSymUsesVisible identifies whether all of the symbol uses of symbols within op are visible.
Applies the transformation script starting at the given transform root operation to the given payload operation. The module containing the transform root as well as the transform options should be provided. The transform operation must implement TransformOpInterface and the module must be a ModuleOp. Returns the status of the application.
Translate operation that satisfies LLVM dialect module requirements into an LLVM IR module living in the given context. This translates operations from any dilalect that has a registered implementation of LLVMTranslationDialectInterface.
Returns
the generated LLVM IR Module from the translated MLIR module, it is owned by the caller.
Prints a location by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Creates an instance of UniformQuantizedPerAxisType with the given parameters in the same context as storageType and returns it. scales and zeroPoints point to nDims number of elements. The instance is owned by the context.
Creates an instance of UniformQuantizedType with the given parameters in the same context as storageType and returns it. The instance is owned by the context.
Unlike the typed accessors below, constructs the attribute with a raw data buffer and no type/alignment checking. Use a more strongly typed accessor if possible. If dataIsMutable is false, then an immutable AsmResourceBlob will be created and that passed data contents will be treated as const. If the deleter is non NULL, then it will be called when the data buffer can no longer be accessed (passing userData to it).
Prints a value by sending chunks of the string representation and forwarding userData tocallback`. Note that the callback may be called several times with consecutive chunks of the string.
Replace all uses of 'of' value with the 'with' value, updating anything in the IR that uses 'of' to use the other value instead. When this returns there are zero uses of 'of'.
Creates a vector type of the shape identified by its rank and dimensions, with the given element type in the same context as the element type. The type is owned by the context.
Creates a scalable vector type with the shape identified by its rank and dimensions. A subset of dimensions may be marked as scalable via the corresponding flag list, which is expected to have as many entries as the rank of the vector. The vector is created in the same context as the element type.
',3))]),t[3185]||(t[3185]=e("h1",{id:"Other-Functions",tabindex:"-1"},[s("Other Functions "),e("a",{class:"header-anchor",href:"#Other-Functions","aria-label":'Permalink to "Other Functions {#Other-Functions}"'},"")],-1))])}const Gk=n(d,[["render",Tk]]);export{Bk as __pageData,Gk as default};
diff --git a/previews/PR363/assets/api_ops.md.D2WkrkEn.js b/previews/PR363/assets/api_ops.md.D2WkrkEn.js
new file mode 100644
index 00000000..a972477f
--- /dev/null
+++ b/previews/PR363/assets/api_ops.md.D2WkrkEn.js
@@ -0,0 +1,15 @@
+import{_ as n,c as l,j as i,a,G as e,a2 as h,B as p,o as k}from"./chunks/framework.2yyKLD8d.js";const C=JSON.parse('{"title":"Reactant.Ops API","description":"","frontmatter":{},"headers":[],"relativePath":"api/ops.md","filePath":"api/ops.md","lastUpdated":null}'),r={name:"api/ops.md"},d={class:"jldocstring custom-block"};function E(o,s,g,c,y,F){const t=p("Badge");return k(),l("div",null,[s[3]||(s[3]=i("h1",{id:"Reactant.Ops-API",tabindex:"-1"},[i("code",null,"Reactant.Ops"),a(" API "),i("a",{class:"header-anchor",href:"#Reactant.Ops-API","aria-label":'Permalink to "`Reactant.Ops` API {#Reactant.Ops-API}"'},"")],-1)),s[4]||(s[4]=i("p",null,[i("code",null,"Reactant.Ops"),a(" module provides a high-level API to construct MLIR operations without having to directly interact with the different dialects.")],-1)),s[5]||(s[5]=i("p",null,[a("Currently we haven't documented all the functions in "),i("code",null,"Reactant.Ops"),a(".")],-1)),i("details",d,[i("summary",null,[s[0]||(s[0]=i("a",{id:"Reactant.Ops.hlo_call-Tuple{Any, Vararg{Any}}",href:"#Reactant.Ops.hlo_call-Tuple{Any, Vararg{Any}}"},[i("span",{class:"jlbinding"},"Reactant.Ops.hlo_call")],-1)),s[1]||(s[1]=a()),e(t,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),s[2]||(s[2]=h(`
Given a MLIR module given as a string, calls the function identified by the func_name keyword parameter (default "main") with the provided arguments and return a tuple for each result of the call.
`,4))])])}const f=n(r,[["render",E]]);export{C as __pageData,f as default};
diff --git a/previews/PR363/assets/api_ops.md.D2WkrkEn.lean.js b/previews/PR363/assets/api_ops.md.D2WkrkEn.lean.js
new file mode 100644
index 00000000..a972477f
--- /dev/null
+++ b/previews/PR363/assets/api_ops.md.D2WkrkEn.lean.js
@@ -0,0 +1,15 @@
+import{_ as n,c as l,j as i,a,G as e,a2 as h,B as p,o as k}from"./chunks/framework.2yyKLD8d.js";const C=JSON.parse('{"title":"Reactant.Ops API","description":"","frontmatter":{},"headers":[],"relativePath":"api/ops.md","filePath":"api/ops.md","lastUpdated":null}'),r={name:"api/ops.md"},d={class:"jldocstring custom-block"};function E(o,s,g,c,y,F){const t=p("Badge");return k(),l("div",null,[s[3]||(s[3]=i("h1",{id:"Reactant.Ops-API",tabindex:"-1"},[i("code",null,"Reactant.Ops"),a(" API "),i("a",{class:"header-anchor",href:"#Reactant.Ops-API","aria-label":'Permalink to "`Reactant.Ops` API {#Reactant.Ops-API}"'},"")],-1)),s[4]||(s[4]=i("p",null,[i("code",null,"Reactant.Ops"),a(" module provides a high-level API to construct MLIR operations without having to directly interact with the different dialects.")],-1)),s[5]||(s[5]=i("p",null,[a("Currently we haven't documented all the functions in "),i("code",null,"Reactant.Ops"),a(".")],-1)),i("details",d,[i("summary",null,[s[0]||(s[0]=i("a",{id:"Reactant.Ops.hlo_call-Tuple{Any, Vararg{Any}}",href:"#Reactant.Ops.hlo_call-Tuple{Any, Vararg{Any}}"},[i("span",{class:"jlbinding"},"Reactant.Ops.hlo_call")],-1)),s[1]||(s[1]=a()),e(t,{type:"info",class:"jlObjectType jlMethod",text:"Method"})]),s[2]||(s[2]=h(`
Given a MLIR module given as a string, calls the function identified by the func_name keyword parameter (default "main") with the provided arguments and return a tuple for each result of the call.
Within each process group in the process grid, concatenates the values of the operand tensor from each process along all_gather_dim and produces a result tensor.
Within each process group in the process grid, applies a reduction function computation to the values of the operand tensor from each process and produces a result tensor.
Within each process group in the process grid, splits the values of the operand tensor along split_dimension into parts, scatters the split parts between the processes, concatenates the scattered parts along concat_dimension and produces a result tensor.
Computes gradients of several inputs of BatchNormTrainingOp backpropagating from grad_output, and produces grad_operand, grad_scale and grad_offset tensors.
Computes mean and variance across batch and spatial dimensions and normalizes the operand tensor, for each feature in the feature_index dimension and produces output, batch_mean and batch_var tensors.
Performs a bitcast operation on operand tensor and produces a result tensor where the bits of the entire operand tensor are reinterpreted using the type of the result tensor.
Within each process group in the process grid, send the value of the operand tensor from the source process to the target processes and produce a result tensor.
Within each process group in the process grid, sends the value of the operand tensor from the source process to the target process and produces a result tensor.
Encapsulates an operation made up (composed) of other StableHLO operations, taking inputs and composite_attributes and producing results. The semantics of the op are implemented by the decomposition attribute. The composite op can be replaced with its decomposition without changing program semantics. In cases where inlining the decomposition does not provide the same op semantics, prefer using custom_call.
The version field (defaults to 0) is used to denote when a composite's semantics change.
This operation is functionally identical to broadcast_in_dim op, but the result shape is specified dynamically via output_dimensions.
It also accepts optional attributes to express static knowledge about the expanding behavior of dimensions. If not specified, all dimensions are assumed to be possibly expanding. The sets of dimensions that are known to be expanding and the set of dimensions that are known to be non-expanding must be disjoint and they must be a subset of the operand's dimensions.
Ensures that the operations that produce the operand are executed before any operations that depend on the result and prevents compiler transformations from moving operations across the barrier. Other than that, the operation is an identity, i.e. result = operand.
Performs element-wise conversion of operand to another floating-point type that uses exponent_bits and mantissa_bits and back to the original floating-point type and produces an output tensor.
Within each process group in the process grid, performs reduction, using computations, over the values of the operand tensor from each process, splits the reduction result along scatter_dimension into parts, and scatters the split parts between the processes to produce the result.
Returns an output filled with uniform random data and an updated output state output_state given an initial state initial_state using the pseudorandom number generator algorithm rng_algorithm.
Performs element-wise rounding towards the nearest integer, breaking ties towards the even integer, on the operand tensor and produces a result tensor.
Performs element-wise reciprocal square root operation on operand tensor and produces a result tensor, implementing the rSqrt operation from the IEEE-754 specification.
Produces results tensors which are equal to inputs tensors except that several slices specified by scatter_indices are updated with the values updates using update_computation.
Scatters the values from the source tensor using scatter based on the outcome of reduce_window of the input tensor using select and produces a result tensor.
%result = stablehlo.slice %operand [1:3, 4:8:2]\n : (tensor<3x8xi64>) -> tensor<2x2xi64>\n\n// Same in generic form: the `1:3` above is mapped to the first entry in\n// `start_indices` and `limit_indices`, while `strides` is implicitly 1.\n// The `4:8:2` above is parsed into the second entry of `start_indices`,\n// `limit_indices` and `strides` respectively.\n%result = "stablehlo.slice" (%operand) {\n start_indices = array<i64: 1, 4>,\n limit_indices = array<i64: 3, 8>,\n strides = array<i64: 1, 2>\n} : (tensor<3x8xi64>) -> tensor<2x2xi64>
Sorts a variadic number of tensors in inputs together, according to a custom comparator, along the given dimension and produces a variadic number of tensors as results.
The batch_dims attribute specifies the number of major batch dimensions (0 or more) that act like a multidimensional loop over both the operand and the index.
Performs element-wise conversion of quantized tensor operand to a floating-point tensor result according to the quantization parameters defined by the operand type.
Performs element-wise conversion of floating-point tensor or quantized tensor operand to a quantized tensor result according to the quantization parameters defined by the result type.
Within each process group in the process grid, concatenates the values of the operand tensor from each process along all_gather_dim and produces a result tensor.
Within each process group in the process grid, applies a reduction function computation to the values of the operand tensor from each process and produces a result tensor.
Within each process group in the process grid, splits the values of the operand tensor along split_dimension into parts, scatters the split parts between the processes, concatenates the scattered parts along concat_dimension and produces a result tensor.
Computes gradients of several inputs of BatchNormTrainingOp backpropagating from grad_output, and produces grad_operand, grad_scale and grad_offset tensors.
Computes mean and variance across batch and spatial dimensions and normalizes the operand tensor, for each feature in the feature_index dimension and produces output, batch_mean and batch_var tensors.
Performs a bitcast operation on operand tensor and produces a result tensor where the bits of the entire operand tensor are reinterpreted using the type of the result tensor.
Within each process group in the process grid, send the value of the operand tensor from the source process to the target processes and produce a result tensor.
Within each process group in the process grid, sends the value of the operand tensor from the source process to the target process and produces a result tensor.
Encapsulates an operation made up (composed) of other StableHLO operations, taking inputs and composite_attributes and producing results. The semantics of the op are implemented by the decomposition attribute. The composite op can be replaced with its decomposition without changing program semantics. In cases where inlining the decomposition does not provide the same op semantics, prefer using custom_call.
The version field (defaults to 0) is used to denote when a composite's semantics change.
This operation is functionally identical to broadcast_in_dim op, but the result shape is specified dynamically via output_dimensions.
It also accepts optional attributes to express static knowledge about the expanding behavior of dimensions. If not specified, all dimensions are assumed to be possibly expanding. The sets of dimensions that are known to be expanding and the set of dimensions that are known to be non-expanding must be disjoint and they must be a subset of the operand's dimensions.
Ensures that the operations that produce the operand are executed before any operations that depend on the result and prevents compiler transformations from moving operations across the barrier. Other than that, the operation is an identity, i.e. result = operand.
Performs element-wise conversion of operand to another floating-point type that uses exponent_bits and mantissa_bits and back to the original floating-point type and produces an output tensor.
Within each process group in the process grid, performs reduction, using computations, over the values of the operand tensor from each process, splits the reduction result along scatter_dimension into parts, and scatters the split parts between the processes to produce the result.
Returns an output filled with uniform random data and an updated output state output_state given an initial state initial_state using the pseudorandom number generator algorithm rng_algorithm.
Performs element-wise rounding towards the nearest integer, breaking ties towards the even integer, on the operand tensor and produces a result tensor.
Performs element-wise reciprocal square root operation on operand tensor and produces a result tensor, implementing the rSqrt operation from the IEEE-754 specification.
Produces results tensors which are equal to inputs tensors except that several slices specified by scatter_indices are updated with the values updates using update_computation.
Scatters the values from the source tensor using scatter based on the outcome of reduce_window of the input tensor using select and produces a result tensor.
%result = stablehlo.slice %operand [1:3, 4:8:2]\n : (tensor<3x8xi64>) -> tensor<2x2xi64>\n\n// Same in generic form: the `1:3` above is mapped to the first entry in\n// `start_indices` and `limit_indices`, while `strides` is implicitly 1.\n// The `4:8:2` above is parsed into the second entry of `start_indices`,\n// `limit_indices` and `strides` respectively.\n%result = "stablehlo.slice" (%operand) {\n start_indices = array<i64: 1, 4>,\n limit_indices = array<i64: 3, 8>,\n strides = array<i64: 1, 2>\n} : (tensor<3x8xi64>) -> tensor<2x2xi64>
Sorts a variadic number of tensors in inputs together, according to a custom comparator, along the given dimension and produces a variadic number of tensors as results.
The batch_dims attribute specifies the number of major batch dimensions (0 or more) that act like a multidimensional loop over both the operand and the index.
Performs element-wise conversion of quantized tensor operand to a floating-point tensor result according to the quantization parameters defined by the operand type.
Performs element-wise conversion of floating-point tensor or quantized tensor operand to a quantized tensor result according to the quantization parameters defined by the result type.
',4))}const Ua=$(Fa,[["render",Ra]]),ja=b({__name:"VersionPicker",props:{screenMenu:{type:Boolean}},setup(o){const e=T([]),t=T("Versions"),s=T(!1);Te();const n=()=>typeof window<"u"&&(window.location.hostname==="localhost"||window.location.hostname==="127.0.0.1"),r=()=>{if(typeof window>"u")return"";const{origin:c,pathname:h}=window.location;if(c.includes("github.io")){const y=h.split("/").filter(Boolean),m=y.length>0?`/${y[0]}/`:"/";return`${c}${m}`}else return c},l=()=>new Promise(c=>{if(n()){c(!1);return}const h=setInterval(()=>{window.DOC_VERSIONS&&window.DOCUMENTER_CURRENT_VERSION&&(clearInterval(h),c(!0))},100);setTimeout(()=>{clearInterval(h),c(!1)},5e3)});return F(async()=>{if(!(typeof window>"u")){try{if(n()){const c=["dev"];e.value=c.map(h=>({text:h,link:"/"})),t.value="dev"}else{const c=await l(),h=P(()=>r());if(c&&window.DOC_VERSIONS&&window.DOCUMENTER_CURRENT_VERSION)e.value=window.DOC_VERSIONS.map(y=>({text:y,link:`${h.value}/${y}/`})),t.value=window.DOCUMENTER_CURRENT_VERSION;else{const y=["dev"];e.value=y.map(m=>({text:m,link:`${h.value}/${m}/`})),t.value="dev"}}}catch(c){console.warn("Error loading versions:",c);const h=["dev"],y=P(()=>r());e.value=h.map(m=>({text:m,link:`${y.value}/${m}/`})),t.value="dev"}s.value=!0}}),(c,h)=>s.value?(a(),d(M,{key:0},[!c.screenMenu&&e.value.length>0?(a(),k(Fe,{key:0,item:{text:t.value,items:e.value},class:"VPVersionPicker"},null,8,["item"])):c.screenMenu&&e.value.length>0?(a(),k(Ue,{key:1,text:t.value,items:e.value,class:"VPVersionPicker"},null,8,["text","items"])):_("",!0)],64)):_("",!0)}}),Ga=$(ja,[["__scopeId","data-v-d483b3a6"]]),za=o=>{if(typeof document>"u")return{stabilizeScrollPosition:n=>async(...r)=>n(...r)};const e=document.documentElement;return{stabilizeScrollPosition:s=>async(...n)=>{const r=s(...n),l=o.value;if(!l)return r;const p=l.offsetTop-e.scrollTop;return await he(),e.scrollTop=l.offsetTop-p,r}}},je="vitepress:tabSharedState",J=typeof localStorage<"u"?localStorage:null,Ge="vitepress:tabsSharedState",Ka=()=>{const o=J==null?void 0:J.getItem(Ge);if(o)try{return JSON.parse(o)}catch{}return{}},qa=o=>{J&&J.setItem(Ge,JSON.stringify(o))},Wa=o=>{const e=ot({});D(()=>e.content,(t,s)=>{t&&s&&qa(t)},{deep:!0}),o.provide(je,e)},Ja=(o,e)=>{const t=W(je);if(!t)throw new Error("[vitepress-plugin-tabs] TabsSharedState should be injected");F(()=>{t.content||(t.content=Ka())});const s=T(),n=P({get(){var c;const l=e.value,p=o.value;if(l){const h=(c=t.content)==null?void 0:c[l];if(h&&p.includes(h))return h}else{const h=s.value;if(h)return h}return p[0]},set(l){const p=e.value;p?t.content&&(t.content[p]=l):s.value=l}});return{selected:n,select:l=>{n.value=l}}};let Le=0;const Ya=()=>(Le++,""+Le);function Qa(){const o=Ae();return P(()=>{var s;const t=(s=o.default)==null?void 0:s.call(o);return t?t.filter(n=>typeof n.type=="object"&&"__name"in n.type&&n.type.__name==="PluginTabsTab"&&n.props).map(n=>{var r;return(r=n.props)==null?void 0:r.label}):[]})}const ze="vitepress:tabSingleState",Xa=o=>{me(ze,o)},Za=()=>{const o=W(ze);if(!o)throw new Error("[vitepress-plugin-tabs] TabsSingleState should be injected");return o},xa={class:"plugin-tabs"},er=["id","aria-selected","aria-controls","tabindex","onClick"],tr=b({__name:"PluginTabs",props:{sharedStateKey:{}},setup(o){const e=o,t=Qa(),{selected:s,select:n}=Ja(t,at(e,"sharedStateKey")),r=T(),{stabilizeScrollPosition:l}=za(r),p=l(n),c=T([]),h=m=>{var C;const S=t.value.indexOf(s.value);let V;m.key==="ArrowLeft"?V=S>=1?S-1:t.value.length-1:m.key==="ArrowRight"&&(V=S(a(),d("div",xa,[v("div",{ref_key:"tablist",ref:r,class:"plugin-tabs--tab-list",role:"tablist",onKeydown:h},[(a(!0),d(M,null,E(i(t),V=>(a(),d("button",{id:`tab-${V}-${i(y)}`,ref_for:!0,ref_key:"buttonRefs",ref:c,key:V,role:"tab",class:"plugin-tabs--tab","aria-selected":V===i(s),"aria-controls":`panel-${V}-${i(y)}`,tabindex:V===i(s)?0:-1,onClick:()=>i(p)(V)},N(V),9,er))),128))],544),u(m.$slots,"default")]))}}),nr=["id","aria-labelledby"],sr=b({__name:"PluginTabsTab",props:{label:{}},setup(o){const{uid:e,selected:t}=Za();return(s,n)=>i(t)===s.label?(a(),d("div",{key:0,id:`panel-${s.label}-${i(e)}`,class:"plugin-tabs--content",role:"tabpanel",tabindex:"0","aria-labelledby":`tab-${s.label}-${i(e)}`},[u(s.$slots,"default",{},void 0,!0)],8,nr)):_("",!0)}}),or=$(sr,[["__scopeId","data-v-9b0d03d2"]]),ar=o=>{Wa(o),o.component("PluginTabs",tr),o.component("PluginTabsTab",or)},ir={extends:Ve,Layout(){return Pe(Ve.Layout,null,{"aside-ads-before":()=>Pe(Ua)})},enhanceApp({app:o}){ar(o),o.component("VersionPicker",Ga)}};export{ir as R,yo as c,L as u};
diff --git a/previews/PR363/assets/index.md.BxToCgvT.js b/previews/PR363/assets/index.md.BxToCgvT.js
new file mode 100644
index 00000000..771faf9d
--- /dev/null
+++ b/previews/PR363/assets/index.md.BxToCgvT.js
@@ -0,0 +1,6 @@
+import{_ as a,c as s,a2 as t,o as e}from"./chunks/framework.2yyKLD8d.js";const r=JSON.parse('{"title":"","description":"","frontmatter":{"layout":"home","hero":{"name":"Reactant.jl Docs","text":"Optimizing Julia Functions with MLIR","tagline":"Optimize Julia Functions With MLIR and XLA for High-Performance Execution on CPU, GPU, TPU and more.","actions":[{"theme":"brand","text":"Tutorials","link":"/tutorials"},{"theme":"alt","text":"API Reference 📚","link":"/api/api"},{"theme":"alt","text":"View on GitHub","link":"https://github.com/EnzymeAD/Reactant.jl"}],"image":{"src":"/logo.svg","alt":"Reactant.jl"}},"features":[{"icon":"🚀","title":"Fast & Device Agnostic","details":"Effortlessly execute your code on CPU, GPU, and TPU with MLIR and XLA.","link":"/introduction"},{"icon":"∂","title":"Built-In MLIR AD","details":"Leverage Enzyme-Powered Automatic Differentiation to Differentiate MLIR Functions","link":"/introduction"},{"icon":"🧩","title":"Composable","details":"Executes and optimizes generic Julia code without requiring special rewriting","link":"/introduction"},{"icon":"⚡","title":"Compiler Optimizations","details":"Fancy MLIR Optimizations seamlessly optimize your Julia code","link":"/introduction"}]},"headers":[],"relativePath":"index.md","filePath":"index.md","lastUpdated":null}'),n={name:"index.md"};function l(h,i,p,k,d,o){return e(),s("div",null,i[0]||(i[0]=[t(`
Its easy to install Reactant.jl. Since Reactant.jl is registered in the Julia General registry, you can simply run the following command in the Julia REPL:
julia
julia> using Pkg
+julia> Pkg.add("Reactant")
If you want to use the latest unreleased version of Reactant.jl, you can run the following command:
julia
julia> using Pkg
+julia> Pkg.add(url="https://github.com/EnzymeAD/Reactant.jl")
using Reactant
+Reactant.set_default_backend("cpu")
julia
using Reactant
+Reactant.set_default_backend("gpu")
julia
using Reactant
+Reactant.set_default_backend("tpu")
`,7)]))}const u=a(n,[["render",l]]);export{r as __pageData,u as default};
diff --git a/previews/PR363/assets/index.md.BxToCgvT.lean.js b/previews/PR363/assets/index.md.BxToCgvT.lean.js
new file mode 100644
index 00000000..771faf9d
--- /dev/null
+++ b/previews/PR363/assets/index.md.BxToCgvT.lean.js
@@ -0,0 +1,6 @@
+import{_ as a,c as s,a2 as t,o as e}from"./chunks/framework.2yyKLD8d.js";const r=JSON.parse('{"title":"","description":"","frontmatter":{"layout":"home","hero":{"name":"Reactant.jl Docs","text":"Optimizing Julia Functions with MLIR","tagline":"Optimize Julia Functions With MLIR and XLA for High-Performance Execution on CPU, GPU, TPU and more.","actions":[{"theme":"brand","text":"Tutorials","link":"/tutorials"},{"theme":"alt","text":"API Reference 📚","link":"/api/api"},{"theme":"alt","text":"View on GitHub","link":"https://github.com/EnzymeAD/Reactant.jl"}],"image":{"src":"/logo.svg","alt":"Reactant.jl"}},"features":[{"icon":"🚀","title":"Fast & Device Agnostic","details":"Effortlessly execute your code on CPU, GPU, and TPU with MLIR and XLA.","link":"/introduction"},{"icon":"∂","title":"Built-In MLIR AD","details":"Leverage Enzyme-Powered Automatic Differentiation to Differentiate MLIR Functions","link":"/introduction"},{"icon":"🧩","title":"Composable","details":"Executes and optimizes generic Julia code without requiring special rewriting","link":"/introduction"},{"icon":"⚡","title":"Compiler Optimizations","details":"Fancy MLIR Optimizations seamlessly optimize your Julia code","link":"/introduction"}]},"headers":[],"relativePath":"index.md","filePath":"index.md","lastUpdated":null}'),n={name:"index.md"};function l(h,i,p,k,d,o){return e(),s("div",null,i[0]||(i[0]=[t(`
Its easy to install Reactant.jl. Since Reactant.jl is registered in the Julia General registry, you can simply run the following command in the Julia REPL:
julia
julia> using Pkg
+julia> Pkg.add("Reactant")
If you want to use the latest unreleased version of Reactant.jl, you can run the following command:
julia
julia> using Pkg
+julia> Pkg.add(url="https://github.com/EnzymeAD/Reactant.jl")
Install Julia v1.10 or above. Reactant.jl is available through the Julia package manager. You can enter it by pressing ] in the REPL and then typing add Reactant. Alternatively, you can also do
Reactant provides two new array types at its core, a ConcreteRArray and a TracedRArray. A ConcreteRArray is an underlying buffer to whatever device data you wish to store and can be created by converting from a regular Julia Array.
julia
using Reactant
+
+julia_data = ones(2, 10)
+reactant_data = Reactant.ConcreteRArray(julia_data)
To compile programs using ConcreteRArray's, one uses the compile function, like as follows:
julia
input1 = Reactant.ConcreteRArray(ones(10))
+input2 = Reactant.ConcreteRArray(ones(10))
+
+function sinsum_add(x, y)
+ return sum(sin.(x) .+ y)
+end
+
+f = @compile sinsum_add(input1,input2)
+
+# one can now run the program
+f(input1, input2)
ConcreteRNumber{Float64}(18.414709848078964)
`,14)]))}const c=i(e,[["render",l]]);export{E as __pageData,c as default};
diff --git a/previews/PR363/assets/introduction_index.md.DwLVSwK5.lean.js b/previews/PR363/assets/introduction_index.md.DwLVSwK5.lean.js
new file mode 100644
index 00000000..8eb23cf0
--- /dev/null
+++ b/previews/PR363/assets/introduction_index.md.DwLVSwK5.lean.js
@@ -0,0 +1,24 @@
+import{_ as i,c as a,a2 as n,o as t}from"./chunks/framework.2yyKLD8d.js";const E=JSON.parse('{"title":"Getting Started","description":"","frontmatter":{},"headers":[],"relativePath":"introduction/index.md","filePath":"introduction/index.md","lastUpdated":null}'),e={name:"introduction/index.md"};function l(p,s,h,k,r,d){return t(),a("div",null,s[0]||(s[0]=[n(`
Install Julia v1.10 or above. Reactant.jl is available through the Julia package manager. You can enter it by pressing ] in the REPL and then typing add Reactant. Alternatively, you can also do
Reactant provides two new array types at its core, a ConcreteRArray and a TracedRArray. A ConcreteRArray is an underlying buffer to whatever device data you wish to store and can be created by converting from a regular Julia Array.
julia
using Reactant
+
+julia_data = ones(2, 10)
+reactant_data = Reactant.ConcreteRArray(julia_data)
Its easy to install Reactant.jl. Since Reactant.jl is registered in the Julia General registry, you can simply run the following command in the Julia REPL:
julia
julia> using Pkg
+julia> Pkg.add("Reactant")
If you want to use the latest unreleased version of Reactant.jl, you can run the following command:
julia
julia> using Pkg
+julia> Pkg.add(url="https://github.com/EnzymeAD/Reactant.jl")
Install Julia v1.10 or above. Reactant.jl is available through the Julia package manager. You can enter it by pressing ] in the REPL and then typing add Reactant. Alternatively, you can also do
Reactant provides two new array types at its core, a ConcreteRArray and a TracedRArray. A ConcreteRArray is an underlying buffer to whatever device data you wish to store and can be created by converting from a regular Julia Array.
julia
using Reactant
+
+julia_data = ones(2, 10)
+reactant_data = Reactant.ConcreteRArray(julia_data)
