<!-- Autogenerated by mlir-tblgen; don't manually edit -->
# 'hfusion' Dialect

Hybrid Fusion (HFusion) dialect.

## Operations

### `hfusion.arange` (hfusion::ArangeOp)

Differs from the classical definition of arange slightly with the addition
of an offset(default is 0) and ability to be multi-dimensional (which also
introduces strides). The offset and stride definition are similar to that of
the memref descriptor.

Given a 3D arange op, the value at each position will be:
`arange[i, j, k] = offset + stride[0] * i + stride[1] * j + stride[2] * k`

Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `offset` | index
| `strides` | variadic of index
| `init` | shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensor` | shaped of any type values

### `hfusion.assert` (hfusion::AssertOp)

_Device-side assert for debugging_

Syntax:

```mlir
operation ::= `hfusion.assert` $msg attr-dict $cond `:` type($cond)
```

`hfusion.assert` takes a literal string `msg` and an argument of
scalar or tensor that should be asserted.

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>msg</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `cond` | integer or ranked tensor of any type values

### `hfusion.atomic_cas` (hfusion::AtomicCasOp)

_Atomic Compare-And-Swap (CAS) Op_

Syntax:

```mlir
operation ::= `hfusion.atomic_cas` attr-dict `ins` `(` $input `:` type($input) `)`
              `outs` `(` $dst `:` type($dst) `)`
              (`->` type($output)^)?
```

Compare-And-Swap (CAS) is an atomic operation that consists of three operands:
Memory location (V), Expected old value (A), New value (B).
The semantics of the operation are: the value of V is updated to B,
only if the value of memory location V is equal to the expected old value A. 
The operation returns the original value of V regardless of whether it is updated or not.

Constraints:

  1. The input and output must have the same rank
     and the same element type.

Arguments:

  * `src0`: expected old value
  * `src1`: new value
  * `dst`: memory location in GM

Examples:

```mlir
hfusion.atomic_cas ins(%src0, %src1 : memref<?xf32>, memref<?xf32>) outs(%dst : memref<?xf32>)
%result = hfusion.atomic_cas ins(%src0, %src1 : tensor<?xf32>, tensor<?xf32>) outs(%dst : tensor<?xf32>) -> tensor<?xf32>
```

Traits: `SameOperandsAndResultRank`

Interfaces: `MemoryEffectOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | variadic of Tensor or Memref
| `dst` | Tensor or Memref

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | Tensor or Memref

### `hfusion.atomic_rmw` (hfusion::AtomicRMWOp)

_Atomic RMW Op_

Syntax:

```mlir
operation ::= `hfusion.atomic_rmw` attr-dict `ins` `(` $input `:` type($input) `)`
              `outs` `(` $dst `:` type($dst) `)`
              `atomic_kind` `=` $atomic_kind
              (`->` type($output)^)?
```

Atomic RMW is an atomic operation that consists of three steps:

1. Read the current value of the specified memory address
2. Perform action depending on atomic_kind attr
3. Return the old value read previously
The whole process is atomic, that is, it will not be interrupted by other threads during the operation.

Constraints:

  1. The input memref and output memref must have the same rank
    and the same element type.

Arguments:

  * `src`: new value
  * `dst`: memory location in GM

Examples:

```mlir
hfusion.atomic_rmw ins(%src : memref<?xf32>) outs(%dst : memref<?xf32>) atomic_kind = <add>
%result = hfusion.atomic_rmw ins(%src : tensor<?xf32>) outs(%dst : tensor<?xf32>) atomic_kind = <or> -> tensor<?xf32>
```

Traits: `SameOperandsAndResultRank`

Interfaces: `MemoryEffectOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>atomic_kind</code></td><td>::mlir::hfusion::AtomicKindAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10</summary>{{% markdown %}}Enum cases:
* none (`NONE`)
* add (`ADD`)
* max (`MAX`)
* min (`MIN`)
* and (`AND`)
* or (`OR`)
* xor (`XOR`)
* cas (`CAS`)
* xchg (`XCHG`)
* umax (`UMAX`)
* umin (`UMIN`){{% /markdown %}}</details></td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | Tensor or Memref
| `dst` | Tensor or Memref

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | Tensor or Memref

### `hfusion.atomic_xchg` (hfusion::AtomicXchgOp)

_Atomic Exchange Op_

Syntax:

```mlir
operation ::= `hfusion.atomic_xchg` attr-dict `ins` `(` $input `:` type($input) `)`
              `outs` `(` $dst `:` type($dst) `)`
              (`mask` `(` $mask^ `:` type($mask) `)`)?
              (`->` type($output)^)?
```

Atomic exchange is an atomic operation that consists of three steps:

1. Read the current value of the specified memory address
2. Write the new value to the memory address
3. Return the old value read previously
The whole process is atomic, that is, it will not be interrupted by other threads during the operation.

Constraints:

  1. The input memref and output memref must have the same rank
    and the same element type.

Arguments:

  * `src`: new value
  * `dst`: memory location in GM
  * `mask` : mask the element

Examples:

```mlir
hfusion.atomic_xchg ins(%src : memref<?xf32>) outs(%dst : memref<?xf32>) mask(%m : memref<?xi1>)
%result = hfusion.atomic_xchg ins(%src : tensor<?xf32>) outs(%dst : tensor<?xf32>) mask(%m : memref<?xi1>) -> tensor<?xf32>
```

Traits: `SameOperandsAndResultRank`

Interfaces: `MemoryEffectOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | Tensor or Memref
| `dst` | Tensor or Memref
| `mask` | Tensor or Memref

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | Tensor or Memref

### `hfusion.barrier` (hfusion::BarrierOp)

_Synchronizes all pipelines of a core._

Syntax:

```mlir
operation ::= `hfusion.barrier` attr-dict
```

The "barrier" op synchronizes all pipelines of a core.

### `hfusion.bitcast` (hfusion::BitcastOp)

_Applies the bitcast function elementwise._

Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.cast` (hfusion::CastOp)

_Applies the cast function elementwise._

Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>round_mode</code></td><td>::mlir::hfusion::RoundModeAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6</summary>{{% markdown %}}
    - RINT: round to nearest, tie to even (c language rint)
    - ROUND: round to nearest, tie away from zero (c language round)
    - FLOOR: round to minus infinity (c language floor)
    - CEIL: round to positive infinity (c language ceil)
    - TRUNC: round to zero (c language trunc)
    - ODD: round to odd (Von Neumann rounding)
  {{% /markdown %}}</details></td></tr>
<tr><td><code>enable_overflow</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>cast</code></td><td>::mlir::hfusion::TypeFnAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2</summary>{{% markdown %}}Enum cases:
* cast_signed (`cast_signed`)
* cast_unsigned (`cast_unsigned`)
* bitcast (`bitcast`){{% /markdown %}}</details></td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.compare` (hfusion::CompareOp)

_Applies the compare function fun elementwise._

No Numeric casting is performed on the input operand.
Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>compare_fn</code></td><td>::mlir::hfusion::CompareFnAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9</summary>{{% markdown %}}Enum cases:
* veq (`veq`)
* vne (`vne`)
* vle (`vle`)
* vlt (`vlt`)
* vge (`vge`)
* vgt (`vgt`)
* vule (`vule`)
* vult (`vult`)
* vuge (`vuge`)
* vugt (`vugt`){{% /markdown %}}</details></td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.cumprod` (hfusion::CumprodOp)

_Calculate cumulative product on a certain dim of the input tensor_

Syntax:

```mlir
operation ::= `hfusion.cumprod` $input attr-dict `:` type($input) `cum_dims` `=` $cum_dims `reverse` `=` $reverse `->` type($output)
```

The cumsum operation calculates cumulative product on a certain dim of
the input tensor. The reverse argument specifies the direction of the cumulative product.
Currently only support one cumulation dim.

Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultRank`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cum_dims</code></td><td>::mlir::DenseI64ArrayAttr</td><td>i64 dense array attribute should be in increasing order</td></tr>
<tr><td><code>reverse</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float or 8-bit signless integer or 16-bit signless integer or 32-bit signless integer or 64-bit signless integer values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float or 8-bit signless integer or 16-bit signless integer or 32-bit signless integer or 64-bit signless integer values

### `hfusion.cumsum` (hfusion::CumsumOp)

_Calculate cumulative sum on a certain dim of the input tensor_

Syntax:

```mlir
operation ::= `hfusion.cumsum` $input attr-dict `:` type($input) `cum_dims` `=` $cum_dims `reverse` `=` $reverse `->` type($output)
```

The cumsum operation calculates cumulative sum on a certain dim of the input
tensor.  The reverse argument specifies the direction of the cumulative sum.
Currently only support one cumulation dim.

Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultRank`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>cum_dims</code></td><td>::mlir::DenseI64ArrayAttr</td><td>i64 dense array attribute should be in increasing order</td></tr>
<tr><td><code>reverse</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float or 8-bit signless integer or 16-bit signless integer or 32-bit signless integer or 64-bit signless integer values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float or 8-bit signless integer or 16-bit signless integer or 32-bit signless integer or 64-bit signless integer values

### `hfusion.deinterleave` (hfusion::DeinterleaveOp)

_Constructs two tensors by deinterleaving an input tensor_

Syntax:

```mlir
operation ::= `hfusion.deinterleave` $input custom<HFusionDeinterleave>($channelIndex) attr-dict `:` type($input) `->` type($output)
```

The deinterleave operation constructs two tensors from a single input
tensor. The first result tensor contains the elements from even indexes, and
the second contains elements from odd indexes(both indexes mean the last
dimension index). And it constrains the last dimension size of input must be
multiples of 2.

The channelIndex attribute controls the output behavior:

* -1: Output all channels (returns two tensors, one with even indexes and one with odd indexes)
* 0: Output only channel 0 (even indexes)
* 1: Output only channel 1 (odd indexes)

Traits: `AlwaysSpeculatableImplTrait`, `Commutative`, `SameOperandsAndResultRank`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`, `ReifyRankedShapedTypeOpInterface`

Effects: `MemoryEffects::Effect{}`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>channelIndex</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | variadic of ranked tensor of any type values

### `hfusion.elemwise_binary` (hfusion::ElemwiseBinaryOp)

_Applies the binary function fun elementwise._

Numeric casting is performed on the input operand, promoting it to the same
data type as the accumulator/output.
Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>fun</code></td><td>::mlir::hfusion::BinaryFnAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17</summary>{{% markdown %}}Enum cases:
* vor (`vor`)
* vand (`vand`)
* vxor (`vxor`)
* minf (`minf`)
* maxf (`maxf`)
* powf (`powf`)
* mod (`mod`)
* modui (`modui`)
* shli (`shli`)
* shrsi (`shrsi`)
* shrui (`shrui`)
* ldexp (`ldexp`)
* ceildivsi (`ceildivsi`)
* ceildivui (`ceildivui`)
* floordivsi (`floordivsi`)
* powi (`powi`)
* minnumf (`minnumf`)
* maxnumf (`maxnumf`){{% /markdown %}}</details></td></tr>
<tr><td><code>cast</code></td><td>::mlir::hfusion::TypeFnAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2</summary>{{% markdown %}}Enum cases:
* cast_signed (`cast_signed`)
* cast_unsigned (`cast_unsigned`)
* bitcast (`bitcast`){{% /markdown %}}</details></td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.elemwise_unary` (hfusion::ElemwiseUnaryOp)

_Applies the unary function fun elementwise._

Numeric casting is performed on the input operand, promoting it to the same
data type as the accumulator/output.
Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>fun</code></td><td>::mlir::hfusion::UnaryFnAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17</summary>{{% markdown %}}Enum cases:
* relu (`relu`)
* sqrt (`sqrt`)
* rsqrt (`rsqrt`)
* rec (`rec`)
* vnot (`vnot`)
* tanh (`tanh`)
* sin (`sin`)
* cos (`cos`)
* atan (`atan`)
* tan (`tan`)
* absi (`absi`)
* erf (`erf`)
* log2 (`log2`)
* log10 (`log10`)
* log1p (`log1p`)
* exp2 (`exp2`)
* expm1 (`expm1`)
* ilogb (`ilogb`){{% /markdown %}}</details></td></tr>
<tr><td><code>cast</code></td><td>::mlir::hfusion::TypeFnAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2</summary>{{% markdown %}}Enum cases:
* cast_signed (`cast_signed`)
* cast_unsigned (`cast_unsigned`)
* bitcast (`bitcast`){{% /markdown %}}</details></td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.flip` (hfusion::FlipOp)

_Flips a tensor x along the dimension dim._

Syntax:

```mlir
operation ::= `hfusion.flip` $input attr-dict `:` type($input)
              `flip_axis` `=` $flip_axis
              `->` type($output)
```

Flips a tensor x along the dimension dim.
currently only final dimension supported.

Traits: `AlwaysSpeculatableImplTrait`, `Commutative`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>flip_axis</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of any type values

### `hfusion.gather` (hfusion::GatherOp)

Gathers one axis of the src tensor into a different with the same shape in
all but the gather axis. Corresponds to triton.language.gather.

Given src:tensor<16x16> and index:tensor<16x4> with axis = 1, the op is
equivalent to:

```text
for i in 0 to 16 {
  for j in 0 to 4 {       // Can be tiled without consequence
    for k in 0 to 16 {    // Cannot be tiled without result potentially
                          //   becoming partial, define as gather axis
      output[i][j] = (index[i][j] == k) ? src[i][k] : output[i][j];
    }
  }
}
```

Traits: `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `BiShengIRAggregatedOpInterface`, `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>axis</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `src` | shaped of any type values
| `index` | shaped of any type values
| `init` | shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result` | variadic of tensor of any type values

### `hfusion.group_matmul` (hfusion::GroupMatmulOp)

_Performs grouped matrix multiplications between expert weights and token embeddings.
For each expert, multiplies its weight matrix with its assigned tokens._

Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.histogram` (hfusion::HistogramOp)

_Compute histogram of an integer tensor with optional mask_

Syntax:

```mlir
operation ::= `hfusion.histogram` $input `,` $num_bins (`,` $mask^)? attr-dict `:` type($input) (`,` type($mask)^)? `->` type($output)
```

For each element of the input tensor, increment the corresponding bin in
the output histogram. The number of bins is a compile-time constant
(`I64Attr`) and the output must be a 1D tensor of length = num_bins.
If a mask tensor is provided, only elements with mask[i] = true are counted.

Interfaces: `BiShengIRAggregatedOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>num_bins</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of 8-bit signless integer or 16-bit signless integer or 32-bit signless integer or 64-bit signless integer values
| `mask` | ranked tensor of 1-bit signless integer values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 32-bit signless integer or 64-bit signless integer values

### `hfusion.interleave` (hfusion::InterleaveOp)

_Constructs one tensor by interleaving n input tensors.
      Only support n = 2 now._

Syntax:

```mlir
operation ::= `hfusion.interleave` $input attr-dict `:` type($input) `->` type($output)
```

Interleaves the values of n tensors along their last dimension.
N tensors must have the same shape.
Input tensors and output tensor must have same rank.

Traits: `AlwaysSpeculatableImplTrait`, `Commutative`, `SameOperandsAndResultRank`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`, `ReifyRankedShapedTypeOpInterface`

Effects: `MemoryEffects::Effect{}`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | variadic of ranked tensor of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of any type values

### `hfusion.isfinite` (hfusion::IsFiniteOp)

_Calculates whether elements of a float type tensor is finite._

Syntax:

```mlir
operation ::= `hfusion.isfinite` $input attr-dict `:` type($input) `->` type($output)
```

Calculates whether elements of a float type tensor is finite
(i.e., not NaN or infinity).

Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultRank`

Interfaces: `BiShengIRAggregatedOpInterface`, `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 1-bit signless integer values

### `hfusion.isinf` (hfusion::IsInfOp)

_Calculates whether elements of a float type tensor equal positive
    infinity or negative infinity._

Syntax:

```mlir
operation ::= `hfusion.isinf` $input attr-dict `:` type($input) `->` type($output)
```

Calculates whether elements of a float type tensor is infinity.
Both positive infinity and negative infinity work.

Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultRank`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 1-bit signless integer values

### `hfusion.isnan` (hfusion::IsNanOp)

_Calculates whether elements of a float type tensor is NAN._

Syntax:

```mlir
operation ::= `hfusion.isnan` $input attr-dict `:` type($input) `->` type($output)
```

Calculates whether elements of a float type tensor is NAN.

Traits: `AlwaysSpeculatableImplTrait`, `SameOperandsAndResultRank`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `input` | ranked tensor of bfloat16 type or 16-bit float or 32-bit float values

#### Results

| Result | Description |
| :----: | ----------- |
| `output` | ranked tensor of 1-bit signless integer values

### `hfusion.load` (hfusion::LoadOp)

_Loads the tensor elementwise._

No numeric casting is performed on the input operand.
Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.mulext` (hfusion::MulExtOp)

_Extended signed integer multiplication operation_

Syntax:

```mlir
operation ::= `hfusion.mulext` $lhs `,` $rhs attr-dict `:` type($lhs)
```

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 with the same operands.

Traits: `AlwaysSpeculatableImplTrait`, `Commutative`

Interfaces: `ConditionallySpeculatable`, `InferTypeOpInterface`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `lhs` | signless-integer-like
| `rhs` | signless-integer-like

#### Results

| Result | Description |
| :----: | ----------- |
| `low` | signless-integer-like
| `high` | signless-integer-like

### `hfusion.print` (hfusion::PrintOp)

_Device-side print for debugging_

Syntax:

```mlir
operation ::= `hfusion.print` $prefix attr-dict $arg `:` type($arg)
```

`hfusion.print` takes a literal string `prefix` and an argument of
scalar or tensor that should be printed. The optional arg `hex` configs
if printing in hex format.

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>prefix</code></td><td>::mlir::StringAttr</td><td>string attribute</td></tr>
<tr><td><code>hex</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `arg` | integer or floating-point or ranked tensor of any type values

### `hfusion.reduce_with_index` (hfusion::ReduceWithIndexOp)

_Max/min reduce with index._

Using max/min to perform a reduce operation on an AnyShaped.
It supports two modes: (1) take input (AnyShaped) and
index (AnyShaped), produce resulting input and index;
(2) take input, produce resulting input and index.
Currently only one reduction dimension is supported.
`tie_break_left` shows whether find the leftmost index or
the rightmost index. 

Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>reduce_kind</code></td><td>::mlir::hfusion::ReduceWithIndexKindAttr</td><td><details><summary></summary>{{% markdown %}}
    The kind of reduce with index.
  {{% /markdown %}}</details></td></tr>
<tr><td><code>tie_break_left</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>dimensions</code></td><td>::mlir::DenseI64ArrayAttr</td><td>i64 dense array attribute should be in increasing order</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of shaped of any type values
| `inits` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result` | variadic of tensor of any type values

### `hfusion.select` (hfusion::SelectOp)

_Chooses one value based on a binary condition supplied as its first operand._

Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.sort` (hfusion::SortOp)

_Sort Op_

Syntax:

```mlir
operation ::= `hfusion.sort` attr-dict `ins` `(` $src `:` type($src) `)`
              `descending` `=` $descending
              `sort_axis` `=` $sort_axis
              (`->` type($result)^)?
```

Sort the sorting axis of `src` in ascending or descending order, and output
the sorted value and the index corresponding to the value.

Constraints:

  1. The input vector and output vector must have the same rank.
  2. Currently only tail axis sorting is supported.

Arguments:

  * `src`: the tensor/memref from which to be sorted
  * `dst_value`: the tensor/memref to store the sorted value
  * `dst_index`: the tensor/memref to store the index corresponding to dst_value
  * `descending`: determines whether to sort in ascending or descending
                  order. The default is false, which means ascending order
  * `sort_axis`: Axis to be sorted

Examples:

  ```mlir
  %result = hfusion.sort ins(%src : tensor<?xf32>) descending = true sort_axis = 0 -> tensor<?xf32>
  ```

Traits: `SameOperandsAndResultRank`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>descending</code></td><td>::mlir::BoolAttr</td><td>bool attribute</td></tr>
<tr><td><code>sort_axis</code></td><td>::mlir::IntegerAttr</td><td>64-bit signless integer attribute</td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `src` | Tensor or Memref

#### Results

| Result | Description |
| :----: | ----------- |
| `result` | variadic of ranked tensor of any type values

### `hfusion.store` (hfusion::StoreOp)

_Stores the tensor elementwise._

No numeric casting is performed on the input operand.
Traits: `AttrSizedOperandSegments`, `SingleBlockImplicitTerminator<mlir::linalg::YieldOp>`, `SingleBlock`

Interfaces: `DestinationStyleOpInterface`, `LinalgStructuredInterface`, `MemoryEffectOpInterface`, `ReifyRankedShapedTypeOpInterface`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>atomic_kind</code></td><td>::mlir::hfusion::AtomicKindAttr</td><td><details><summary>allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10</summary>{{% markdown %}}Enum cases:
* none (`NONE`)
* add (`ADD`)
* max (`MAX`)
* min (`MIN`)
* and (`AND`)
* or (`OR`)
* xor (`XOR`)
* cas (`CAS`)
* xchg (`XCHG`)
* umax (`UMAX`)
* umin (`UMIN`){{% /markdown %}}</details></td></tr>
</table>

#### Operands

| Operand | Description |
| :-----: | ----------- |
| `inputs` | variadic of any type
| `outputs` | variadic of shaped of any type values

#### Results

| Result | Description |
| :----: | ----------- |
| `result_tensors` | variadic of ranked tensor of any type values

### `hfusion.symbolic_dim` (hfusion::SymbolicDimOp)

_Symbolic dimension reference returning an index_

Syntax:

```mlir
operation ::= `hfusion.symbolic_dim` $symbolName attr-dict `:` type($result)
```

The "symbolic_dim" operation references a symbolic dimension by name
(via a symbol attribute) and returns it as an `index`-typed value.

Traits: `AlwaysSpeculatableImplTrait`

Interfaces: `ConditionallySpeculatable`, `InferTypeOpInterface`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Attributes

<table>
<tr><th>Attribute</th><th>MLIR Type</th><th>Description</th></tr>
<tr><td><code>symbolName</code></td><td>::mlir::SymbolRefAttr</td><td>symbol reference attribute</td></tr>
</table>

#### Results

| Result | Description |
| :----: | ----------- |
| `result` | index

## Attributes

### AtomicKindAttr

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10

Syntax:

```mlir
#hfusion.atomic_kind<
  ::mlir::hfusion::AtomicKind   # value
>
```

Enum cases:

* none (`NONE`)
* add (`ADD`)
* max (`MAX`)
* min (`MIN`)
* and (`AND`)
* or (`OR`)
* xor (`XOR`)
* cas (`CAS`)
* xchg (`XCHG`)
* umax (`UMAX`)
* umin (`UMIN`)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::AtomicKind` | an enum of type AtomicKind |

### BinaryFnAttr

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17

Syntax:

```mlir
#hfusion.binary_fn<
  ::mlir::hfusion::BinaryFn   # value
>
```

Enum cases:

* vor (`vor`)
* vand (`vand`)
* vxor (`vxor`)
* minf (`minf`)
* maxf (`maxf`)
* powf (`powf`)
* mod (`mod`)
* modui (`modui`)
* shli (`shli`)
* shrsi (`shrsi`)
* shrui (`shrui`)
* ldexp (`ldexp`)
* ceildivsi (`ceildivsi`)
* ceildivui (`ceildivui`)
* floordivsi (`floordivsi`)
* powi (`powi`)
* minnumf (`minnumf`)
* maxnumf (`maxnumf`)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::BinaryFn` | an enum of type BinaryFn |

### CompareFnAttr

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9

Syntax:

```mlir
#hfusion.compare_fn<
  ::mlir::hfusion::CompareFn   # value
>
```

Enum cases:

* veq (`veq`)
* vne (`vne`)
* vle (`vle`)
* vlt (`vlt`)
* vge (`vge`)
* vgt (`vgt`)
* vule (`vule`)
* vult (`vult`)
* vuge (`vuge`)
* vugt (`vugt`)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::CompareFn` | an enum of type CompareFn |

### BindSubBlockAttr

Syntax: `#hfusion.bind_sub_block`

Specific operations for bind sub block.

### FusionKindAttr

Syntax:

```mlir
#hfusion.fusion_kind<
  ::mlir::hfusion::FusionKind   # fusion_kind
>
```

HFusion fused kernel kind.

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| fusion_kind | `::mlir::hfusion::FusionKind` | an enum of type FusionKind |

### InsertSliceSourceIndexAttr

Syntax: `#hfusion.insert_slice_source_index`

Specifies which operand is insert_slice source in concat op

### MultiBufferAttr

Syntax: `#hfusion.multi_buffer`

HFusion multi buffer attribute for target op.

### ReduceComposeAttr

Syntax: `#hfusion.reduce_composed`

HFusion reduced composed.

### ReduceWithIndexKindAttr

Syntax:

```mlir
#hfusion.reduce_with_index_kind<
  ::mlir::hfusion::ReduceWithIndexKind   # reduce_with_index_kind
>
```

The kind of reduce with index.

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| reduce_with_index_kind | `::mlir::hfusion::ReduceWithIndexKind` | an enum of type ReduceWithIndexKind |

### ReturnOperandNumAttr

Syntax: `#hfusion.return_operand_num`

Specifies which operand this corresponds to in the function return

### StrideAlignDimsAttr

Syntax: `#hfusion.stride_align_dims`

HFusion stride align dims.

### StrideAlignValueInByteAttr

Syntax: `#hfusion.stride_align_value_in_byte`

HFusion stride align value in byte.

### RoundModeAttr

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6

Syntax:

```mlir
#hfusion.round_mode<
  ::mlir::hfusion::RoundMode   # value
>
```

- RINT: round to nearest, tie to even (c language rint)
- ROUND: round to nearest, tie away from zero (c language round)
- FLOOR: round to minus infinity (c language floor)
- CEIL: round to positive infinity (c language ceil)
- TRUNC: round to zero (c language trunc)
- ODD: round to odd (Von Neumann rounding)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::RoundMode` | an enum of type RoundMode |

### TernaryFnAttr

allowed 32-bit signless integer cases: 0

Syntax:

```mlir
#hfusion.ternary_fn<
  ::mlir::hfusion::TernaryFn   # value
>
```

Enum cases:

* select (`select`)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::TernaryFn` | an enum of type TernaryFn |

### TypeFnAttr

allowed 32-bit signless integer cases: 0, 1, 2

Syntax:

```mlir
#hfusion.type_fn<
  ::mlir::hfusion::TypeFn   # value
>
```

Enum cases:

* cast_signed (`cast_signed`)
* cast_unsigned (`cast_unsigned`)
* bitcast (`bitcast`)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::TypeFn` | an enum of type TypeFn |

### UnaryFnAttr

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17

Syntax:

```mlir
#hfusion.unary_fn<
  ::mlir::hfusion::UnaryFn   # value
>
```

Enum cases:

* relu (`relu`)
* sqrt (`sqrt`)
* rsqrt (`rsqrt`)
* rec (`rec`)
* vnot (`vnot`)
* tanh (`tanh`)
* sin (`sin`)
* cos (`cos`)
* atan (`atan`)
* tan (`tan`)
* absi (`absi`)
* erf (`erf`)
* log2 (`log2`)
* log10 (`log10`)
* log1p (`log1p`)
* exp2 (`exp2`)
* expm1 (`expm1`)
* ilogb (`ilogb`)

#### Parameters

| Parameter | C++ type | Description |
| :-------: | :-------: | ----------- |
| value | `::mlir::hfusion::UnaryFn` | an enum of type UnaryFn |

## Enums

### AtomicKind

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| NONE | `0` | none |
| ADD | `1` | add |
| MAX | `2` | max |
| MIN | `3` | min |
| AND | `4` | and |
| OR | `5` | or |
| XOR | `6` | xor |
| CAS | `7` | cas |
| XCHG | `8` | xchg |
| UMAX | `9` | umax |
| UMIN | `10` | umin |

### BinaryFn

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| vor | `0` | vor |
| vand | `1` | vand |
| vxor | `2` | vxor |
| minf | `3` | minf |
| maxf | `4` | maxf |
| powf | `5` | powf |
| mod | `6` | mod |
| modui | `7` | modui |
| shli | `8` | shli |
| shrsi | `9` | shrsi |
| shrui | `10` | shrui |
| ldexp | `11` | ldexp |
| ceildivsi | `12` | ceildivsi |
| ceildivui | `13` | ceildivui |
| floordivsi | `14` | floordivsi |
| powi | `15` | powi |
| minnumf | `16` | minnumf |
| maxnumf | `17` | maxnumf |

### CastMode

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| F32TOI8 | `0` | F32TOI8 |
| F32TOI16 | `1` | F32TOI16 |
| F16TOI8 | `2` | F16TOI8 |
| I64TOI32 | `3` | I64TOI32 |
| I64TOI16 | `4` | I64TOI16 |
| I64TOI8 | `5` | I64TOI8 |
| I32TOI16 | `6` | I32TOI16 |
| I32TOI8 | `7` | I32TOI8 |
| I16TOI8 | `8` | I16TOI8 |

### CompareFn

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| veq | `0` | veq |
| vne | `1` | vne |
| vle | `2` | vle |
| vlt | `3` | vlt |
| vge | `4` | vge |
| vgt | `5` | vgt |
| vule | `6` | vule |
| vult | `7` | vult |
| vuge | `8` | vuge |
| vugt | `9` | vugt |

### FlattenMode

HFusion flatten mode

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| Greedy | `1` | Greedy |
| Tidy | `2` | Tidy |

### FusionKind

HFusion fused kernel kind

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| PureElemwise | `1` | PURE_ELEMWISE |
| AnyPB | `2` | ANY_PB |
| LastAxisPBR | `3` | LAST_AXIS_PBR |
| AnyPBR | `4` | ANY_PBR |
| SingleCube | `5` | SINGLE_CUBE |
| ShallowCV | `6` | SHALLOW_CV |
| ShallowVV | `7` | SHALLOW_VV |
| MixCV | `8` | MIX_CV |
| MixC2 | `9` | MIX_C2 |
| Unknown | `10` | UNKNOWN |

### OutputMode

HFusion Output mode

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| Multiple | `1` | Multiple |
| Single | `2` | Single |
| SingleAggressive | `3` | SingleAggressive |

### CumOpType

HFusion cumulative operation type

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| UNDEFINED | `0` | undefined |
| CUMSUM | `1` | cumsum |
| CUMPROD | `2` | cumprod |

### MmMapMode

allowed 32-bit signless integer cases: 0, 1

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| CoreOp | `0` | core_op |
| MacroInstr | `1` | macro_instr |

### ReduceWithIndexKind

allowed 32-bit signless integer cases: 0, 1, 2, 3

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| MIN | `0` | min |
| MAX | `1` | max |
| MINUI | `2` | minui |
| MAXUI | `3` | maxui |

### RoundMode

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| RINT | `0` | rint |
| ROUND | `1` | round |
| FLOOR | `2` | floor |
| CEIL | `3` | ceil |
| TRUNC | `4` | trunc |
| ODD | `5` | odd |
| TRUNCWITHOVERFLOW | `6` | truncwithoverflow |

### TaylorMode

allowed 32-bit signless integer cases: 0, 1

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| SIN | `0` | sin |
| ATAN | `1` | atan |

### TernaryFn

allowed 32-bit signless integer cases: 0

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| select | `0` | select |

### TypeFn

allowed 32-bit signless integer cases: 0, 1, 2

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| cast_signed | `0` | cast_signed |
| cast_unsigned | `1` | cast_unsigned |
| bitcast | `2` | bitcast |

### UnaryFn

allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17

#### Cases

| Symbol | Value | String |
| :----: | :---: | ------ |
| relu | `0` | relu |
| sqrt | `1` | sqrt |
| rsqrt | `2` | rsqrt |
| rec | `3` | rec |
| vnot | `4` | vnot |
| tanh | `5` | tanh |
| sin | `6` | sin |
| cos | `7` | cos |
| atan | `8` | atan |
| tan | `9` | tan |
| absi | `10` | absi |
| erf | `11` | erf |
| log2 | `12` | log2 |
| log10 | `13` | log10 |
| log1p | `14` | log1p |
| exp2 | `15` | exp2 |
| expm1 | `16` | expm1 |
| ilogb | `17` | ilogb |
Attribute	MLIR Type	Description
`round_mode`	::mlir::hfusion::RoundModeAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6 {{% markdown %}} - RINT: round to nearest, tie to even (c language rint) - ROUND: round to nearest, tie away from zero (c language round) - FLOOR: round to minus infinity (c language floor) - CEIL: round to positive infinity (c language ceil) - TRUNC: round to zero (c language trunc) - ODD: round to odd (Von Neumann rounding) {{% /markdown %}}
`enable_overflow`	::mlir::BoolAttr	bool attribute
`cast`	::mlir::hfusion::TypeFnAttr	allowed 32-bit signless integer cases: 0, 1, 2 {{% markdown %}}Enum cases: * cast_signed (`cast_signed`) * cast_unsigned (`cast_unsigned`) * bitcast (`bitcast`){{% /markdown %}}
Attribute	MLIR Type	Description
`cum_dims`	::mlir::DenseI64ArrayAttr	i64 dense array attribute should be in increasing order
`reverse`	::mlir::BoolAttr	bool attribute
Attribute	MLIR Type	Description
`fun`	::mlir::hfusion::BinaryFnAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17 {{% markdown %}}Enum cases: * vor (`vor`) * vand (`vand`) * vxor (`vxor`) * minf (`minf`) * maxf (`maxf`) * powf (`powf`) * mod (`mod`) * modui (`modui`) * shli (`shli`) * shrsi (`shrsi`) * shrui (`shrui`) * ldexp (`ldexp`) * ceildivsi (`ceildivsi`) * ceildivui (`ceildivui`) * floordivsi (`floordivsi`) * powi (`powi`) * minnumf (`minnumf`) * maxnumf (`maxnumf`){{% /markdown %}}
`cast`	::mlir::hfusion::TypeFnAttr	allowed 32-bit signless integer cases: 0, 1, 2 {{% markdown %}}Enum cases: * cast_signed (`cast_signed`) * cast_unsigned (`cast_unsigned`) * bitcast (`bitcast`){{% /markdown %}}
Attribute	MLIR Type	Description
`fun`	::mlir::hfusion::UnaryFnAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16, 17 {{% markdown %}}Enum cases: * relu (`relu`) * sqrt (`sqrt`) * rsqrt (`rsqrt`) * rec (`rec`) * vnot (`vnot`) * tanh (`tanh`) * sin (`sin`) * cos (`cos`) * atan (`atan`) * tan (`tan`) * absi (`absi`) * erf (`erf`) * log2 (`log2`) * log10 (`log10`) * log1p (`log1p`) * exp2 (`exp2`) * expm1 (`expm1`) * ilogb (`ilogb`){{% /markdown %}}
`cast`	::mlir::hfusion::TypeFnAttr	allowed 32-bit signless integer cases: 0, 1, 2 {{% markdown %}}Enum cases: * cast_signed (`cast_signed`) * cast_unsigned (`cast_unsigned`) * bitcast (`bitcast`){{% /markdown %}}
Attribute	MLIR Type	Description
`prefix`	::mlir::StringAttr	string attribute
`hex`	::mlir::BoolAttr	bool attribute
Attribute	MLIR Type	Description
`reduce_kind`	::mlir::hfusion::ReduceWithIndexKindAttr	{{% markdown %}} The kind of reduce with index. {{% /markdown %}}
`tie_break_left`	::mlir::BoolAttr	bool attribute
`dimensions`	::mlir::DenseI64ArrayAttr	i64 dense array attribute should be in increasing order
Attribute	MLIR Type	Description
`descending`	::mlir::BoolAttr	bool attribute
`sort_axis`	::mlir::IntegerAttr	64-bit signless integer attribute