‘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:

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¶

Attribute	MLIR Type	Description
`msg`	::mlir::StringAttr	string attribute

Operands¶

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

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

Atomic Compare-And-Swap (CAS) Op

Syntax:

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:

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:

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:

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:

Read the current value of the specified memory address
Perform action depending on atomic_kind attr
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:

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:

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¶

Attribute	MLIR Type	Description
`atomic_kind`	::mlir::hfusion::AtomicKindAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10 {{% 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 %}}

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:

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:

Read the current value of the specified memory address
Write the new value to the memory address
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:

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:

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:

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¶

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 %}}

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¶

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

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:

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¶

Attribute	MLIR Type	Description
`cum_dims`	::mlir::DenseI64ArrayAttr	i64 dense array attribute should be in increasing order
`reverse`	::mlir::BoolAttr	bool attribute

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:

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¶

Attribute	MLIR Type	Description
`cum_dims`	::mlir::DenseI64ArrayAttr	i64 dense array attribute should be in increasing order
`reverse`	::mlir::BoolAttr	bool attribute

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:

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¶

Attribute	MLIR Type	Description
`channelIndex`	::mlir::IntegerAttr	64-bit signless integer attribute

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¶

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 %}}

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.

Interfaces: DestinationStyleOpInterface, LinalgStructuredInterface, MemoryEffectOpInterface, ReifyRankedShapedTypeOpInterface

Attributes¶

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 %}}

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:

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¶

Attribute	MLIR Type	Description
`flip_axis`	::mlir::IntegerAttr	64-bit signless integer attribute

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:

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¶

Attribute	MLIR Type	Description
`axis`	::mlir::IntegerAttr	64-bit signless integer attribute

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:

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¶

Attribute	MLIR Type	Description
`num_bins`	::mlir::IntegerAttr	64-bit signless integer attribute

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:

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:

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:

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:

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:

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:

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¶

Attribute	MLIR Type	Description
`prefix`	::mlir::StringAttr	string attribute
`hex`	::mlir::BoolAttr	bool attribute

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¶

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

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:

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:

The input vector and output vector must have the same rank.
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:

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

Traits: SameOperandsAndResultRank

Attributes¶

Attribute	MLIR Type	Description
`descending`	::mlir::BoolAttr	bool attribute
`sort_axis`	::mlir::IntegerAttr	64-bit signless integer attribute

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¶

Attribute	MLIR Type	Description
`atomic_kind`	::mlir::hfusion::AtomicKindAttr	allowed 32-bit signless integer cases: 0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10 {{% 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 %}}

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:

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¶

Attribute	MLIR Type	Description
`symbolName`	::mlir::SymbolRefAttr	symbol reference attribute

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:

#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:

#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:

#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:

#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:

#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:

#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:

#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:

#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:

#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

‘hfusion’ Dialect¶

Operations¶

hfusion.arange (hfusion::ArangeOp)¶

Operands¶

Results¶

hfusion.assert (hfusion::AssertOp)¶

Attributes¶

Operands¶

hfusion.atomic_cas (hfusion::AtomicCasOp)¶

Operands¶

Results¶

hfusion.atomic_rmw (hfusion::AtomicRMWOp)¶

Attributes¶

Operands¶

Results¶

hfusion.atomic_xchg (hfusion::AtomicXchgOp)¶

Operands¶

Results¶

hfusion.barrier (hfusion::BarrierOp)¶

hfusion.bitcast (hfusion::BitcastOp)¶

Operands¶

Results¶

hfusion.cast (hfusion::CastOp)¶

Attributes¶

Operands¶

Results¶

hfusion.compare (hfusion::CompareOp)¶

Attributes¶

Operands¶

Results¶

hfusion.cumprod (hfusion::CumprodOp)¶

Attributes¶

Operands¶

Results¶

hfusion.cumsum (hfusion::CumsumOp)¶

Attributes¶

Operands¶

Results¶

hfusion.deinterleave (hfusion::DeinterleaveOp)¶

Attributes¶

Operands¶

Results¶

hfusion.elemwise_binary (hfusion::ElemwiseBinaryOp)¶

Attributes¶

Operands¶

Results¶

hfusion.elemwise_unary (hfusion::ElemwiseUnaryOp)¶

Attributes¶

Operands¶

Results¶

hfusion.flip (hfusion::FlipOp)¶

Attributes¶

Operands¶

Results¶

hfusion.gather (hfusion::GatherOp)¶

Attributes¶

Operands¶

Results¶

hfusion.group_matmul (hfusion::GroupMatmulOp)¶

Operands¶

Results¶

hfusion.histogram (hfusion::HistogramOp)¶

Attributes¶

Operands¶

Results¶

hfusion.interleave (hfusion::InterleaveOp)¶

Operands¶

Results¶

hfusion.isfinite (hfusion::IsFiniteOp)¶

Operands¶

Results¶

hfusion.isinf (hfusion::IsInfOp)¶

Operands¶

Results¶

hfusion.isnan (hfusion::IsNanOp)¶

Operands¶

Results¶

hfusion.load (hfusion::LoadOp)¶

Operands¶

Results¶

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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