「ムロ」方言

運営

mhlo.abs (mhlo::AbsOp)

腹筋手術

構文:

operation ::= `mhlo.abs` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

operandテンソルに対して要素ごとの abs 演算を実行し、 resultテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#abs

例:

%result = mhlo.abs %operand : tensor<3xi32>

特性: AlwaysSpeculatableImplTraitElementwiseSameOperandsAndResultShape

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

オペランド:

オペランド説明
operand 2/4/8/16/32/64 ビットの符号なし整数のランク付きテンソル、または f4E2M1FN 型、f6E2M3FN 型、または f6E3M2FN 型、または f8E3M4 型、または f8E4M3 型、または f8E4M3FN 型、または f8E4M3FNUZ 型、または f8E4M3B11FNUZ 型、または f8E5M2 型、またはf8E5M2FNUZ 型、f8E8M0FNU 型、または 16 ビット float、32 ビット float、64 ビット float、または bfloat16 型、または 32 ビット float または 64 ビット float 要素の複合型、または 2/4/8/16/32 ビットuniform量子化された符号付き整数または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数または2/4/8/16/32 ビット均一量子化符号なし整数値、または軸ごとの 2/4/8/16/32 ビット均一量子化符号なし整数値

結果:

結果説明
result 2/4/8/16/32/64 ビットの符号なし整数のランク付きテンソル、または f4E2M1FN 型、f6E2M3FN 型、または f6E3M2FN 型、または f8E3M4 型、または f8E4M3 型、または f8E4M3FN 型、または f8E4M3FNUZ 型、または f8E4M3B11FNUZ 型、または f8E5M2 型、またはf8E5M2FNUZ 型、f8E8M0FNU 型、または 16 ビット float、または 32 ビット float、または 64 ビット float、または bfloat16 型、または 2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32-軸ごとに均一に量子化されたビット、符号付き整数、または 2/4/8/16/32 ビット均一に量子化符号なし整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値

mhlo.add (mhlo::AddOp)

追加操作

構文:

operation ::= `mhlo.add` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

2 つのテンソルlhsrhsの要素ごとの加算を実行し、 resultテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#add

例:

%result = mhlo.add %lhs, %rhs : tensor<2x2xi32>

特性: AlwaysSpeculatableImplTraitCommutativeCompatibleOperandsAndResultTypeElementwiseSameOperandsAndResultShape

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

オペランド:

オペランド説明
lhs f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8 /16/32/64 ビットの符号なし整数または 32 ビット浮動小数点数または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値
rhs f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8 /16/32/64 ビットの符号なし整数または 32 ビット浮動小数点数または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

結果:

結果説明
result f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8 /16/32/64 ビットの符号なし整数または 32 ビット浮動小数点数または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

mhlo.add_dependency (mhlo::AddDependencyOp)

AddDependency 操作

構文:

operation ::= `mhlo.add_dependency` operands attr-dict `:` functional-type(operands, results)

この操作は XLA コンパイラーにプライベートなものであるため、まだ仕様がありません。

非公式には、この操作にはデータ オペランドとトークンの 2 つのオペランドがあります。演算の出力はデータ オペランドです。 AfterAll とともに使用すると、この操作により、副作用のない操作 (トークン値を生成しない操作) の順序付けが可能になります。

例:

%1 = mhlo.add_dependency %arg0, %0 : (tensor<3x4xf32>, !mhlo.token) -> tensor<3x4xf32>

特性: AlwaysSpeculatableImplTrait

インターフェイス: ConditionallySpeculatableInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

オペランド:

オペランド説明
operand f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8 /16/32/64 ビットの符号なし整数または 32 ビット浮動小数点数または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数値、または 2/4/8/16/32 ビット均一量子化のランク付きテンソル軸の符号付き整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値またはトークン
tokenトークン

結果:

結果説明
output f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8 /16/32/64 ビットの符号なし整数または 32 ビット浮動小数点数または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数値、または 2/4/8/16/32 ビット均一量子化のランク付きテンソル軸の符号付き整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値またはトークン

mhlo.after_all (mhlo::AfterAllOp)

アフターオール操作

構文:

operation ::= `mhlo.after_all` $inputs attr-dict
              `:` custom<VariadicSameOperandsAndResultType>(ref($inputs), type($inputs), type($result))

inputs生成する操作が、 resultに依存する操作の前に実行されるようにします。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all

例:

%result = mhlo.after_all %input0, %input1 : !mhlo.token

特性: AlwaysSpeculatableImplTrait

インターフェイス: ConditionallySpeculatableInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

オペランド:

オペランド説明
inputsトークンの可変個引数

結果:

結果説明
resultトークン

mhlo.all_gather (mhlo::AllGatherOp)

オールギャザー操作

プロセス グリッド内の各プロセス グループ内で、各プロセスからのオペランド テンソルの値をall_gather_dimに沿って連結し、結果テンソルを生成します。 computation operandsのオペランドごとに個別に適用され、オペランドごとに 1 つの結果が生成されます。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_gather

例:

%result = "mhlo.all_gather"(%operand) {
  all_gather_dim = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>,
  // use_global_device_ids = false
} : (tensor<2x2xf32>) -> tensor<2x4xf32>

特性: SameOperandsAndResultElementType

属性:

属性MLIRタイプ説明
all_gather_dim ::mlir::IntegerAttr値が負でない 64 ビットの符号なし整数属性
replica_groups ::mlir::DenseIntElementsAttr 64 ビット符号なし整数要素属性
channel_handle ::mlir::mhlo::ChannelHandleAttr 2 つの 64 ビット整数「ハンドル」と「タイプ」
use_global_device_ids ::mlir::UnitAttrユニット属性

オペランド:

オペランド説明
operands f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

結果:

結果説明
«無名» f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

mhlo.all_reduce (mhlo::AllReduceOp)

AllReduce 操作

プロセス グリッド内の各プロセス グループ内で、各プロセスからのオペランド テンソルの値にリダクション関数のcomputationを適用し、結果テンソルを生成します。 computation operandsのオペランドごとに個別に適用され、オペランドごとに 1 つの結果が生成されます。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce

例:

%result = "mhlo.all_reduce"(%operand) ({
  ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>):
    %0 = mhlo.add %arg1, %arg2 : tensor<f32>
    mhlo.return %0 : tensor<f32>
}) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<4xf32>) -> tensor<4xf32>

特性: InferTensorTypeSingleBlockImplicitTerminator<ReturnOp>SingleBlock

インターフェイス: InferShapedTypeOpInterfaceInferTypeOpInterface

属性:

属性MLIRタイプ説明
replica_groups ::mlir::DenseIntElementsAttr 64 ビット符号なし整数要素属性
channel_handle ::mlir::mhlo::ChannelHandleAttr 2 つの 64 ビット整数「ハンドル」と「タイプ」
use_global_device_ids ::mlir::UnitAttrユニット属性

オペランド:

オペランド説明
operands f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

結果:

結果説明
«無名» f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

mhlo.all_to_all (mhlo::AllToAllOp)

AllToAll 操作

プロセス グリッドの各プロセス グループ内で、 split_dimensionに沿ってoperandテンソルの値を部分に分割し、分割した部分をプロセス間で分散させ、分散した部分をconcat_dimensionに沿って連結して、 resultテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_to_all

例:

%result = "mhlo.all_to_all"(%operand) {
  split_dimension = 1 : i64,
  concat_dimension = 0 : i64,
  split_count = 2 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
} : (tensor<2x4xf32>) -> tensor<4x2xf32>

特性: AlwaysSpeculatableImplTraitInferTensorTypeSameOperandsElementTypeSameOperandsShapeSameVariadicOperandSize

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
split_dimension ::mlir::IntegerAttr値が負でない 64 ビットの符号なし整数属性
concat_dimension ::mlir::IntegerAttr値が負でない 64 ビットの符号なし整数属性
split_count ::mlir::IntegerAttr値が正の 64 ビットの符号なし整数属性
replica_groups ::mlir::DenseIntElementsAttr 64 ビット符号なし整数要素属性
channel_handle ::mlir::mhlo::ChannelHandleAttr 2 つの 64 ビット整数「ハンドル」と「タイプ」

オペランド:

オペランド説明
operand f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

結果:

結果説明
«無名» f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

mhlo.and (mhlo::AndOp)

そして操作

構文:

operation ::= `mhlo.and` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

2 つのテンソルlhsrhsの要素ごとの AND を実行し、 resultテンソルを生成します

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#and

例:

%result = mhlo.and %lhs, %rhs : tensor<2x2xi32>

特性: AlwaysSpeculatableImplTraitCommutativeCompatibleOperandsAndResultTypeElementwiseSameOperandsAndResultShape

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

オペランド:

オペランド説明
lhs pred のランク付きテンソル (ブール値または 1 ビット整数)、2/4/8/16/32/64 ビットの符号なし整数、または 2/4/8/16/32/64 ビットの符号なし整数値
rhs pred のランク付きテンソル (ブール値または 1 ビット整数)、2/4/8/16/32/64 ビットの符号なし整数、または 2/4/8/16/32/64 ビットの符号なし整数値

結果:

結果説明
result f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8 /16/32/64 ビットの符号なし整数または 32 ビット浮動小数点数または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または2/4/8/16/32 ビットの軸ごとに均一に量子化された符号なし整数値

mhlo.async_done (mhlo::AsyncDoneOp)

非同期完了操作

この操作は XLA コンパイラーにプライベートなものであるため、まだ仕様がありません。

非公式には、この操作は非同期計算が終了するまでブロックされます。非同期計算の最終結果を返します。

詳細については、AsyncStart のドキュメントを参照してください。

インターフェイス: InferTypeOpInterface

オペランド:

オペランド説明
bundle f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、または f8E5M2FNUZ 型のランク付けされたテンソルの任意の組み合わせを含む async_bundle またはf8E8M0FNU 型、16 ビット float、32 ビット float、64 ビット float、bfloat16 型、pred (ブール型または 1 ビット整数)、2/4/8/16/32/64 ビット符号なし整数、または 2/ 4/8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点または 64 ビットの複素数型float 要素または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/4/8/16/32 ビット均一量子化符号なし整数または軸ごとの 2/4/8/16/32 ビット均一量子化符号付き整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値またはトークン値

結果:

結果説明
«無名» f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値、トークン、またはランク付けされたテンソルの任意の組み合わせを含むネストされたタプルf4E2M1FNタイプ or f6E2M3FNタイプ or f6E3M2FNタイプ or f8E3M4タイプ or f8E4M3タイプ or f8E4M3FNタイプ or f8E4M3FNUZタイプ or f8E4M3B11FNUZタイプ or f8E5M2タイプ or f8E5M2FNUZタイプ or f8E8M0FNUタイプ or 16ビットfloatタイプ or 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8/16/32/ 64 ビット符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素、または 2/4/8/16/32 ビットを含む複合型均一量子化された符号付き整数または 2/4/8/16/32 ビット均一量子化された符号なし整数値または 2/4/8/16/32 ビットのランク付きテンソル 軸ごとに均一量子化された符号付き整数または 2/4/8/16 /軸ごとに均一に量子化された 32 ビットの符号なし整数値またはトークン値

mhlo.async_start (mhlo::AsyncStartOp)

非同期開始操作

この操作は XLA コンパイラーにプライベートなものであるため、まだ仕様がありません。

非公式には、この操作により非同期計算が開始されます。

これは、非同期待機 (DMA など) とオンスレッド計算の両方を含む関数がある場合に使用されます。たとえば、関数は、計算、DMA、別の計算、2 番目の DMA、および最終計算で構成される場合があります。これは、async_start とそれに続く async_update および async_done として表されます。 async_start はスレッド上で最初の計算を実行し、その後 DMA を開始します。 async_update は、DMA がまだ完了していない場合は完了するまで待機し、関数内の 2 番目の計算を実行して、2 番目の DMA を開始します。最後に、async_done はこの最後の DMA を待機してから、スレッド上で実行する必要がある最後の計算を実行し、その最後の計算の結果を返します。

operands計算に直接渡されますcalled_computation非同期で実行される関数ですexecution_thread実行されるスレッドの名前です。メインスレッドを「メイン」と呼びます。すべてのスレッドには名前があります。

これにより、非同期操作の間に必要なすべての状態が返されます。バッファ割り当て後の戻り値は、入力、結果、および非同期操作によって必要または編集されたスクラッチパッドを保持するために必要なスペースを表します。

属性:

属性MLIRタイプ説明
called_computation ::mlir::FlatSymbolRefAttrフラット シンボル参照属性
execution_thread ::mlir::StringAttr文字列属性

オペランド:

オペランド説明
inputs f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、または f8E8M0FNU のランク付けされたテンソルの可変個数type または 16 ビット float または 32 ビット float または 64 ビット float または bfloat16 type または pred (別名ブール値または 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4 /8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素を含む複合型、または2/4/8/16/32 ビット均一量子化符号付き整数、または 2/4/8/16/32 ビット均一量子化符号なし整数、または 2/4/8/16/32 ビット均一量子化された軸ごとの符号付き整数、または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値、トークン、またはランク付けされたテンソルの任意の組み合わせを含むネストされたタプルf4E2M1FNタイプ or f6E2M3FNタイプ or f6E3M2FNタイプ or f8E3M4タイプ or f8E4M3タイプ or f8E4M3FNタイプ or f8E4M3FNUZタイプ or f8E4M3B11FNUZタイプ or f8E5M2タイプ or f8E5M2FNUZタイプ or f8E8M0FNUタイプ or 16ビットfloatタイプ or 32 ビット float または 64 ビット float または bfloat16 型または pred (別名ブールまたは 1 ビット整数) または 2/4/8/16/32/64 ビット符号なし整数または 2/4/8/16/32/ 64 ビット符号なし整数、または 32 ビット浮動小数点要素または 64 ビット浮動小数点要素、または 2/4/8/16/32 ビットを含む複合型均一量子化された符号付き整数または 2/4/8/16/32 ビット均一量子化された符号なし整数値または 2/4/8/16/32 ビットのランク付きテンソル 軸ごとに均一量子化された符号付き整数または 2/4/8/16 /軸ごとに均一に量子化された 32 ビットの符号なし整数値またはトークン値

結果:

結果説明
«無名» f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、または f8E5M2FNUZ 型のランク付けされたテンソルの任意の組み合わせを含む async_bundle またはf8E8M0FNU 型、16 ビット float、32 ビット float、64 ビット float、bfloat16 型、pred (ブール型または 1 ビット整数)、2/4/8/16/32/64 ビット符号なし整数、または 2/ 4/8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点または 64 ビットの複素数型float 要素または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/4/8/16/32 ビット均一量子化符号なし整数または軸ごとの 2/4/8/16/32 ビット均一量子化符号付き整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値またはトークン値

mhlo.async_update (mhlo::AsyncUpdateOp)

非同期更新操作

この操作は XLA コンパイラーにプライベートなものであるため、まだ仕様がありません。

非公式には、この操作は同期バリアに達するまで非同期計算をブロックします。これにより、操作後にbundleが返されます。

詳細については、AsyncStart のドキュメントを参照してください。

インターフェイス: InferTypeOpInterface

オペランド:

オペランド説明
bundle f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、または f8E5M2FNUZ 型のランク付けされたテンソルの任意の組み合わせを含む async_bundle またはf8E8M0FNU 型、16 ビット float、32 ビット float、64 ビット float、bfloat16 型、pred (ブール型または 1 ビット整数)、2/4/8/16/32/64 ビット符号なし整数、または 2/ 4/8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点または 64 ビットの複素数型float 要素または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/4/8/16/32 ビット均一量子化符号なし整数または軸ごとの 2/4/8/16/32 ビット均一量子化符号付き整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値またはトークン値

結果:

結果説明
«無名» f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、または f8E5M2FNUZ 型のランク付けされたテンソルの任意の組み合わせを含む async_bundle またはf8E8M0FNU 型、16 ビット float、32 ビット float、64 ビット float、bfloat16 型、pred (ブール型または 1 ビット整数)、2/4/8/16/32/64 ビット符号なし整数、または 2/ 4/8/16/32/64 ビットの符号なし整数、または 32 ビット浮動小数点または 64 ビットの複素数型float 要素または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/4/8/16/32 ビット均一量子化符号なし整数または軸ごとの 2/4/8/16/32 ビット均一量子化符号付き整数値または軸ごとに均一に量子化された 2/4/8/16/32 ビットの符号なし整数値またはトークン値

mhlo.atan2 (mhlo::Atan2Op)

Atan2の動作

構文:

operation ::= `mhlo.atan2` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

lhsおよびrhsテンソルに対して要素ごとの atan2 演算を実行し、 resultテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#atan2

例:

%result = mhlo.atan2 %lhs, %rhs : tensor<3xf32>

特性: AlwaysSpeculatableImplTraitCompatibleOperandsAndResultTypeElementwiseSameOperandsAndResultShape

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

オペランド:

オペランド説明
lhs f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型、または 32 ビット float または 64 ビット float 要素を含む複合型、または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/ 4/8/16/32 ビットの均一量子化符号なし整数値
rhs f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型、または 32 ビット float または 64 ビット float 要素を含む複合型、または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/ 4/8/16/32 ビットの均一量子化符号なし整数値

結果:

結果説明
result f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型、または 32 ビット float または 64 ビット float 要素を含む複合型、または 2/4/8/16/32 ビット均一量子化符号付き整数または 2/ 4/8/16/32 ビットの均一量子化符号なし整数値

mhlo.batch_norm_grad (mhlo::BatchNormGradOp)

BatchNormGrad 操作

grad_outputから逆伝播する BatchNormTrainingOp のいくつかの入力の勾配を計算し、 grad_operandgrad_scale 、およびgrad_offsetテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_grad

例:

%grad_operand, %grad_scale, %grad_offset =
"mhlo.batch_norm_grad"(%operand, %scale, %mean, %variance, %grad_output) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>,
    tensor<2x2x2xf32>) -> (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>)

特性: AlwaysSpeculatableImplTraitInferTensorType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
epsilon ::mlir::FloatAttr 32ビット浮動小数点属性
feature_index ::mlir::IntegerAttr値が負でない 64 ビットの符号なし整数属性

オペランド:

オペランド説明
operand f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
scale f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
mean f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型の値
variance f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型の値
grad_output f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型の値

結果:

結果説明
grad_operand f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
grad_scale f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
grad_offset f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値

mhlo.batch_norm_inference (mhlo::BatchNormInferenceOp)

BatchNormInference 操作

feature_index次元を除くすべての次元にわたってoperandテンソルを正規化し、 resultテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_inference

例:

%result = "mhlo.batch_norm_inference"(%operand, %scale, %offset, %mean, %variance) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> tensor<2x2x2xf32>

特性: AlwaysSpeculatableImplTraitInferTensorType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
epsilon ::mlir::FloatAttr 32ビット浮動小数点属性
feature_index ::mlir::IntegerAttr値が負でない 64 ビットの符号なし整数属性

オペランド:

オペランド説明
operand f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
scale f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
offset f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値
mean f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float または 32 ビット float または 64 ビット float または bfloat16 型の値
variance f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型の 1D テンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値

結果:

結果説明
result f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16 ビット float、32 ビット float、64 ビット float、または bfloat16 型の値

mhlo.batch_norm_training (mhlo::BatchNormTrainingOp)

BatchNormTraining オペレーション

バッチ次元と空間次元にわたる平均と分散を計算し、 feature_index次元の各特徴についてoperandテンソルを正規化し、 outputbatch_meanテンソル、 batch_varテンソルを生成します。

参照: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_training

例:

%output, %batch_mean, %batch_var = "mhlo.batch_norm_training"(%operand, %scale, %offset) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>) -> (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>)

特性: AlwaysSpeculatableImplTraitInferTensorType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

エフェクト: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
epsilon ::mlir::FloatAttr 32ビット浮動小数点属性
feature_index ::mlir::IntegerAttr値が負でない 64 ビットの符号なし整数属性

オペランド:

オペランド説明
operand f4E2M1FN 型、f6E2M3FN 型、f6E3M2FN 型、f8E3M4 型、f8E4M3 型、f8E4M3FN 型、f8E4M3FNUZ 型、f8E4M3B11FNUZ 型、f8E5M2 型、f8E5M2FNUZ 型、f8E8M0FNU 型のランク付けされたテンソル、または16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ値
scale F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E3M4タイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E3M4タイプまたはF8E3M4タイプの1Dテンソル16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ値
offset F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E3M4タイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E3M4タイプまたはF8E3M4タイプの1Dテンソル16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ値

結果:

結果説明
output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ値
batch_mean F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E3M4タイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E3M4タイプまたはF8E3M4タイプの1Dテンソル16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ値
batch_var F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E3M4タイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E4M3FNタイプまたはF8E4M3FNタイプまたはF8E3M4タイプまたはF8E3M4タイプの1Dテンソル16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ値

mhlo.bitcast (mhlo :: bitcastop)

ビットキャスト操作

構文:

operation ::= `mhlo.bitcast` operands attr-dict `:` functional-type(operands, results)

この操作はXLAコンパイラのプライベートであるため、まだ仕様がありません。

非公式には、この操作は、要素の物理的な配置が変化しないように入力の形状を変更します。

この操作には、「要素の物理的配置」を理解するためにレイアウト情報が必要であり、MHLOでのレイアウトサポートは現在進行中です。

例:

%0 = mhlo.bitcast %arg0 : (tensor<3x4xf32>) -> tensor<3x4x1xf32>

特性: AlwaysSpeculatableImplTrait

インターフェイス: ConditionallySpeculatableNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.bitcast_convert (mhlo :: bitcastconvertop)

BitCastConvert操作

構文:

operation ::= `mhlo.bitcast_convert` operands attr-dict `:` functional-type(operands, results)

operandテンソルでビットキャスト操作を実行し、 result operandのタイプを使用して再解釈されるresultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#bitcast_convert

例:

%result = mhlo.bitcast_convert %operand : (tensor<2xf32>) -> tensor<2x4xi8>

特性: AlwaysSpeculatableImplTrait

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.broadcast (mhlo :: broadcastop)

ブロードキャスト操作

この操作はStablehloから出ているので、仕様には含まれていません:https: //github.com/openxla/stablehlo/issues/3

非公式には、この操作はXLAの放送と同じことをします: https://www.tensorflow.org/xla/operation_semantics#broadcast

例:

%result = mhlo.broadcast %operand, sizes = [1, 2] : (tensor<3xi32>) -> tensor<1x2x3xi32>

特性: AlwaysSpeculatableImplTrait yimpltrait、 InferTensorTypeSameOperandsAndResultElementType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
broadcast_sizes :: mlir :: denseintelementsatttr 64ビットサインレス整数要素属性

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.broadcast_in_dim (mhlo :: broadcastindimop)

broadcastindim操作

operandテンソルのデータを複製することにより、入力テンソルの寸法および/またはランクを拡張し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#broadcast_in_dim

例:

%result = mhlo.broadcast_in_dim %operand, dims = [2, 1] : (tensor<1x3xi32>) -> tensor<2x3x2xi32>

特性: AlwaysSpeculatableImplTrait yimpltrait、 HLO_CompatibleOperandsAndResultElementType

インターフェイス: ConditionallySpeculatableNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
broadcast_dimensions :: mlir :: denseintelementsatttr 64ビットサインレス整数要素属性

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプタイプまたはF8E5M2タイプタイプまたはF8E5M2タイプタイプまたはF8E5M2タイプタイプまたはF8E5M2M0FNUZタイプまたはF8E5M2FNUZのタイプのF8E4M3型またはF8E4M3の静的な形状のテンソルタイプまたは16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットのサインレス整数または2/4 /8/16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.case (mhlo :: caseop)

ケース操作

indexの値に応じて、 branchesから正確に1つのfunctionを実行することから出力を生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#case

例:

%result0, %result1 = "mhlo.case"(%index) ({
  mhlo.return %result_branch0, %result_branch0 : tensor<2xi64>, tensor<2xi64>
}, {
  mhlo.return %result_branch1, %result_branch1 : tensor<2xi64>, tensor<2xi64>
}) : (tensor<i32>) -> (tensor<2xi64>, tensor<2xi64>)

特性: RecursiveMemoryEffectsSingleBlockImplicitTerminator<ReturnOp>SingleBlock

インターフェイス: InferTypeOpInterface

オペランド:

オペランド説明
index 32ビットのサインレス整数値のテンソル

結果:

結果説明
«無名» F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZのタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZのタイプのF8E5M2NUZタイプまたはF8E5M2FNU5MUZ5MUZのタイプのランク付けされたテンソルのバリエイジックタイプまたは16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットのサインレス整数または2/4 /8/16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized署名整数整数または2/4/8/16/32ビット均一な量子化されていない整数整数値またはランク付けされたテンソル2/4/8/16/32ビットの均一な均一な定量化軸に署名された整数または2/4/8/16/32ビット均一軸は署名されていない整数値またはトークンごとに量子化されています

mhlo.cbrt (mhlo :: cbrtop)

CBRT操作

構文:

operation ::= `mhlo.cbrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

operandテンソルで要素ごとのキュービックルート操作を実行し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#cbrt

例:

%result = mhlo.cbrt %operand : tensor<4xf32>

特性: AlwaysSpeculatableImplTrait yimpltrait、 CompatibleOperandsAndResultTypeElementwiseSameOperandsAndResultShape

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ32ビットフロートまたは64ビットフロート要素、または2/4/8/16/16/32ビットの均一な量子化された標識整数または2/ 4/8/16/32ビット均一な量子化されていない整数整数値

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプ32ビットフロートまたは64ビットフロート要素、または2/4/8/16/16/32ビットの均一な量子化された標識整数または2/ 4/8/16/32ビット均一な量子化されていない整数整数値

mhlo.ceil (mhlo :: ceilop)

天井操作

構文:

operation ::= `mhlo.ceil` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

operandテンソルの要素ごとの天井を実行し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#ceil

例:

%result = mhlo.ceil %operand : tensor<5xf32>

特性: AlwaysSpeculatableImplTrait yimpltrait、 CompatibleOperandsAndResultTypeElementwiseSameOperandsAndResultShape

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたは2/4/8/16/32ビット均一量子化された署名整数または2/4/8/16/16/32ビット均一な量子化されていない整数整数整数整数整数整数整数整数整合値

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたは2/4/8/16/32ビット均一量子化署名整数または2/4/8/16/16/32ビット均一な量子化されていない整数整数整数整数整数整数整数整数整数整数

mhlo.cholesky (Mhlo :: Choleskyop)

コレスキー操作

マトリックスのバッチの軟骨の分解を計算します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#cholesky

例:

%result = mhlo.cholesky %a, lower = true : tensor<3x3xf32>

特性: AlwaysSpeculatableImplTrait yimpltrait、 InferTensorTypeSameOperandsAndResultElementType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
lower :: MLIR :: BOOLATTRブール属性

オペランド:

オペランド説明
a ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたは32ビットフロートまたは64ビットのフロート要素値を備えた複雑なタイプ

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたは32ビットフロートまたは64ビットのフロート要素値を備えた複雑なタイプ

mhlo.clamp (MHLo :: Clampop)

クランプ操作

構文:

operation ::= `mhlo.clamp` $min `,` $operand `,` $max attr-dict
              `:` custom<SameOperandsAndResultType>(type($min), type($operand), type($max), type($result))

operandテンソルのすべての要素を最小値と最大値の間でクランプし、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#clamp

例:

%result = mhlo.clamp %min, %operand, %max : tensor<3xi32>

特性: AlwaysSpeculatableImplTrait speculatableimpltrait、 HLO_BroadcastingElementwiseInferTensorTypeSameOperandsAndResultElementType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

オペランド:

オペランド説明
min ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値
max ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.collective_broadcast (mhlo :: collectivebroadcastop)

CollectiveBroadcast Operation

プロセスグリッド内の各プロセスグループ内で、ソースプロセスからoperandテンソルの値をターゲットプロセスに送信し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_broadcast

例:

%result = "mhlo.collective_broadcast"(%operand) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
} : (tensor<1x2xi64>) -> tensor<1x2xi64>

特性: CompatibleOperandsAndResultType

インターフェイス: InferShapedTypeOpInterfaceInferTypeOpInterface

属性:

属性MLIRタイプ説明
replica_groups :: mlir :: denseintelementsatttr 64ビットサインレス整数要素属性
channel_handle :: MLIR :: MHLO :: ChannelHandleattr 2つの64ビット整数 'ハンドル'と「タイプ」

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.collective_permute (mhlo :: collectivepermuteop)

CollectivePermute操作

プロセスグリッド内の各プロセスグループ内で、ソースプロセスからターゲットプロセスにoperandテンソルの値を送信し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_permute

例:

%result = "mhlo.collective_permute"(%operand) {
  source_target_pairs = dense<[[0, 1], [1, 2]]> : tensor<2x2xi64>,
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
} : (tensor<4x2xf32>) -> tensor<4x2xf32>

特性: AlwaysSpeculatableImplTraitCompatibleOperandsAndResultType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
source_target_pairs :: mlir :: denseintelementsatttr 64ビットサインレス整数要素属性
channel_handle :: MLIR :: MHLO :: ChannelHandleattr 2つの64ビット整数 'ハンドル'と「タイプ」

オペランド:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.compare (mhlo :: compareop)

操作を比較します

構文:

operation ::= `mhlo.compare` $comparison_direction `,` $lhs `,` $rhs (`,` $compare_type^)?
              attr-dict `:` functional-type(operands, results)

comparison_directioncompare_typeに従って、 lhsrhsテンソルの要素ごとの比較を実行し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#compare

例:

%result = mhlo.compare LT, %lhs, %rhs, FLOAT : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>

特性: AlwaysSpeculatableImplTraitElementwiseInferTensorTypeSameOperandsAndResultShapeSameOperandsElementType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
comparison_direction :: mlir :: mhlo :: ComparisondirectionAttr実行する比較操作。
compare_type :: mlir :: mhlo :: compationtypeattr使用する比較タイプ。

オペランド:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» pred(別名ブールまたは1ビット整数)値のランク付けされたテンソル

mhlo.complex (mhlo :: complexop)

複雑な操作

構文:

operation ::= `mhlo.complex` operands attr-dict
              `:` custom<ComplexOpType>(type($lhs), type($rhs), type($result))

実際の値と架空の値、 lhsおよびrhsペアから複雑な値への要素ごとの変換を実行し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#complex

例:

%result = mhlo.complex %lhs, %rhs : tensor<2xcomplex<f32>>

特性: AlwaysSpeculatableImplTrait yimpltrait、 ElementwiseSameOperandsAndResultShapeSameOperandsElementType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

オペランド:

オペランド説明
lhs 32ビットフロートまたは64ビットフロート値のランク付けテンソル
rhs 32ビットフロートまたは64ビットフロート値のランク付けテンソル

結果:

結果説明
result 32ビットフロートまたは64ビットフロート要素の値を持つ複雑なタイプのランク付けされたテンソル

mhlo.composite (mhlo :: compositeop)

複合操作

構文:

operation ::= `mhlo.composite` $name $inputs attr-dict `:` functional-type(operands, results)

他の安定した操作で構成された(構成された)操作をカプセル化し、 inputscomposite_attributesを取得し、 resultsを生成します。 OPのセマンティクスは、 decomposition属性によって実装されます。 composite OPは、プログラムセマンティクスを変更することなく、分解に置き換えることができます。分解が同じOPセマンティクスを提供しない場合、 custom_callの使用を好みます。

versionフィールド(デフォルトは0 )を使用して、コンポジットのセマンティクスが変更されるときに示すために使用されます。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#composite

例:

%results = mhlo.composite "my.op" %arg0, %arg1 {
  decomposition = @my_op,
  composite_attributes = { my_attribute = "my_value" },
  version = 1 : i32
} : (tensor<f32>, tensor<f32>) -> tensor<f32>

インターフェイス: SymbolUserOpInterface

属性:

属性MLIRタイプ説明
name :: mlir :: stringattr文字列属性
composite_attributes :: mlir :: dictionaryattr名前付き属性値の辞書
decomposition :: mlir :: flatsymbolrefattrフラットシンボル参照属性
version :: mlir :: integerattr 32ビットサインレス整数属性

オペランド:

オペランド説明
inputs F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZのタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZのタイプのF8E5M2NUZタイプまたはF8E5M2FNU5MUZ5MUZのタイプのランク付けされたテンソルのバリエイジックタイプまたは16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットのサインレス整数または2/4 /8/16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット軸ごとに量子化された均一な均一な整列整数値またはトークンまたはネストされたタプルのランク付けされたテンソルの組み合わせF4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E3M4タイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプ32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットサインレス整数または2/4/8/16/16/32/ 32ビットフロートまたは64ビットフロート要素、または2/4/8/16/32ビットを備えた64ビットの非署名整数または複雑なタイプ均一な量子化された署名整数または2/4/8/16/32ビット均一均一な量子化されていない整数整数値またはランク付けされたテンソル2/4/8/16/32ビット均一均一標識整数または2/4/8/16 /軸ごとに32ビットの均一な量子化されていない整数整数値またはトークン値

結果:

結果説明
«無名» F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZのタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZのタイプのF8E5M2NUZタイプまたはF8E5M2FNU5MUZ5MUZのタイプのランク付けされたテンソルのバリエイジックタイプまたは16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットのサインレス整数または2/4 /8/16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット軸ごとに量子化された均一な均一な整列整数値またはトークンまたはネストされたタプルのランク付けされたテンソルの組み合わせF4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E3M4タイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプ32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットサインレス整数または2/4/8/16/16/32/ 32ビットフロートまたは64ビットフロート要素、または2/4/8/16/32ビットを備えた64ビットの非署名整数または複雑なタイプ均一な量子化された署名整数または2/4/8/16/32ビット均一均一な量子化されていない整数整数値またはランク付けされたテンソル2/4/8/16/32ビット均一均一標識整数または2/4/8/16 /軸ごとに32ビットの均一な量子化されていない整数整数値またはトークン値

mhlo.concatenate (mhlo :: concatenateop)

連結動作

指定された引数と同じ順序で、 dimension寸法に沿ったinputsの変形数のテンソルを連結し、 resultテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#concatenate

例:

%result = mhlo.concatenate %input0, %input1, dim = 0 : (tensor<3x2xi64>, tensor<1x2xi64>) -> tensor<4x2xi64>

特性: AlwaysSpeculatableImplTraitSameOperandsAndResultElementType

インターフェイス: ConditionallySpeculatableInferShapedTypeOpInterfaceInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
dimension :: mlir :: integerattr値が非陰性である64ビットサインレス整数属性

オペランド:

オペランド説明
val F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZ5MUZのタイプまたはF8E5M2FNUZ5MUZ5MUZ5MUZ5MUZのタイプのF8E5M2NUZタイプまたはF8E5M2FNU5MUZ5MUZのタイプのランク付けされたテンソルのバリエイジックタイプまたは16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPRED(別名ブールまたは1ビット整数)または2/4/8/16/32/64ビットのサインレス整数または2/4 /8/16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

結果:

結果説明
«無名» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16ビットフロートまたは32ビットフロートまたは64ビットフロートまたはBFLOAT16タイプまたはPred(別名ブールまたは1ビット整数)または2/4/8/16/64ビットサインレス整数または2/4/8 /16/32/64ビット32ビットフロートまたは64ビットフロート要素を備えた符号なし整数または複雑なタイプまたは2/4/8/16/32ビット均一均一Quantized Signed Signed Integerまたは2/4/8/16/32ビット均一Quantized Unsigned Unsigned Integerまたは2/4/8/16/16/32ビット均一な均一な均一な統合整数または整数または2/4/8/16/32ビット均一軸ごとに量子化されていない統一された整数整数値

mhlo.constant (mhlo :: constantop)

一定の動作

一定のvalueからoutputテンソルを生成します。

参照:https: //github.com/openxla/stablehlo/blob/main/docs/spec.md#constant

例:

%output = mhlo.constant dense<[[0.0, 1.0], [2.0, 3.0]]> : tensor<2x2xf32>

特性: AlwaysSpeculatableImplTraitConstantLike

インターフェイス: ConditionallySpeculatableInferTypeOpInterfaceNoMemoryEffect (MemoryEffectOpInterface)

効果: MemoryEffects::Effect{}

属性:

属性MLIRタイプ説明
value :: mlir :: elementaTtr定数ベクトル/テンソル属性

結果:

結果説明
output F4E2M1FNタイプまたはF6E2M3FNタイプまたはF6E3M2FNタイプまたはF8E4M3タイプまたはF8E4M3FNタイプまたはF8E4M3FNUZタイプまたはF8E4M3B11FNUZタイプまたはF8E5M2タイプタイプまたはF8E5M2タイプタイプまたはF8E5M2タイプタイプまたはF8E5M2タイプタイプまたはF8E5M2M0FNUZタイプまたはF8E5M2FNU5FNUZタイプまたはF8E5M2FNUZPNU5FNUZタイプまたはF8E5M2FNUZのタイプの静的な形状のテンソル type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.convert (mhlo::ConvertOp)

Convert operation

構文:

operation ::= `mhlo.convert` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs an element-wise conversion from one element type to another on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convert

例:

%result = mhlo.convert %operand : (tensor<3xi32>) -> tensor<3xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.convolution (mhlo::ConvolutionOp)

Convolution operation

構文:

operation ::= `mhlo.convolution` `(`operands`)`
              `dim_numbers` `=` custom<ConvolutionDimensions>($dimension_numbers) `,`
              `window` `=` `{` custom<WindowAttributes>($window_strides, $padding,
              $lhs_dilation, $rhs_dilation,
              $window_reversal) `}`
              attr-dict `:` functional-type(operands, results)

Computes dot products between windows of lhs and slices of rhs and produces result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convolution

例:

%result = "mhlo.convolution"(%lhs, %rhs) {
  window_strides = dense<4> : tensor<2xi64>,
  padding = dense<0> : tensor<2x2xi64>,
  lhs_dilation = dense<2> : tensor<2xi64>,
  rhs_dilation = dense<1> : tensor<2xi64>,
  window_reversal = dense<false> : tensor<2xi1>,
  dimension_numbers = #mhlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi32>, tensor<3x3x1x1xi32>) -> tensor<1x2x2x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
lhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
rhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_reversal ::mlir::DenseElementsAttr constant boolean vector/tensor attribute
dimension_numbers ::mlir::mhlo::ConvDimensionNumbersAttr Structure of dimension information for conv op
feature_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.copy (mhlo::CopyOp)

Copy operation

構文:

operation ::= `mhlo.copy` operands attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation a copy of operand . Depending on the metadata attached to the operation, it can behave quite differently from a no-op.

例:

%0 = mhlo.copy %arg0 : tensor<f32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
cross_program_prefetch_index ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.cosine (mhlo::CosineOp)

Cosine operation

構文:

operation ::= `mhlo.cosine` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise cosine operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cosine

例:

%result = mhlo.cosine %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.count_leading_zeros (mhlo::ClzOp)

Clz operation

構文:

operation ::= `mhlo.count_leading_zeros` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise count of the number of leading zero bits in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#count_leading_zeros

例:

%result = mhlo.count_leading_zeros %operand : tensor<2x2xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.create_token (mhlo::CreateTokenOp)

CreateToken operation

構文:

operation ::= `mhlo.create_token` attr-dict `:` type(results)

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as AfterAllOp with 0 inputs: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all

例:

%output = mhlo.create_token : !mhlo.token

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

結果:

結果説明
outputトークン

mhlo.cross-replica-sum (mhlo::CrossReplicaSumOp)

CrossReplicaSum operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as AllReduceOp with channel_id = 0 , use_global_device_ids = false and computation implementing addition: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce

例:

%result = "mhlo.cross-replica-sum"(%operand) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
} : (tensor<4xf32>) -> tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.custom_call (mhlo::CustomCallOp)

CustomCall operation

構文:

operation ::= `mhlo.custom_call` custom<CustomCallTarget>($call_target_name) `(` $inputs `)`
              attr-dict `:` functional-type(operands, results)

Encapsulates an implementation-defined operation call_target_name that takes inputs and called_computations and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#custom_call

例:

%results = "mhlo.custom_call"(%input0) {
  call_target_name = "foo",
  has_side_effect = false,
  backend_config = "bar",
  api_version = 1 : i32,
  called_computations = [@foo]
} : (tensor<f32>) -> tensor<f32>

A custom call invokes code external to XLA. The `inputs` are passed to the
external code, and the external code is expected to produce a result of the
given type. The exact mechanism is backend-specific. For example, in the CPU
backend, a call instruction is emitted which targets a symbol with the name
`call_target_name`.

If XLA runtime is enabled for a backend, then custom calls use the runtime
custom call calling convention to call into the external functions. This
calling convention defines an ABI for encoding arguments, attributes and
results.

Depending on the API version there are two ways to pass extra bits of static
information to the external function:

1. For `API_VERSION_TYPED_FFI` custom calls `backend_config` must be a
   dictionary attribute, that will be encoded according to the custom call
   calling convention and passed to the external function as the attributes
   argument. External code is expected to use declarative bindings (see
   `xla/runtime/custom_call.h`) to decode them at run time. These custom
   calls are only supported if XLA uses XLA runtime.

2. For previous API versions it is the user responsibility to encode extra
   bits of static information as a string `backend_config` attribute, and
   decode it at run time.

Interfaces: MemoryEffectOpInterface

Attributes:

属性MLIR Type説明
call_target_name ::mlir::StringAttr string attribute
has_side_effect ::mlir::BoolAttr bool attribute
backend_config ::mlir::Attribute string attribute or dictionary of named attribute values
api_version ::mlir::mhlo::CustomCallApiVersionAttr Custom call API version
called_computations ::mlir::ArrayAttr flat symbol ref array attribute
custom_call_schedule ::mlir::mhlo::CustomCallScheduleAttr Specifies the desired schedule for the custom-call.
operand_layouts ::mlir::ArrayAttr Array of layout (1D tensor of index type) attributes
result_layouts ::mlir::ArrayAttr Array of layout (1D tensor of index type) attributes
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of CustomCall

Operands:

オペランド説明
inputs variadic of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

結果:

結果説明
«unnamed» variadic of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.divide (mhlo::DivOp)

Div operation

構文:

operation ::= `mhlo.divide` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise division of dividend lhs and divisor rhs tensors and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#divide

例:

%result = mhlo.divide %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.domain (mhlo::DomainOp)

Domain operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, these operations are used to group instructions with the same DomainMetadata property. ShardingMetadata is the main use case today to group instructions on the same device. Domain instructions provide two major benefits:

  • Prevent unintentionally optimizing instructions across domains.
  • Automatically assign the metadata of the instructions created in the domain. Without domain instructions, each HLO optimization pass would have to check and propagate the metadata, which would be easy to miss and also adds complexity to the compiler. Since domain instructions connect two different domains, each domain instruction is associated with two DomainMetadata -- one on the operand side and one on the user side of the domain.

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
kind ::mlir::mhlo::DomainKindAttr Kind of domain metatdata attached to an HLO domain.
entry_metadata ::mlir::StringAttr string attribute
exit_metadata ::mlir::StringAttr string attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.dot (mhlo::DotOp)

Dot operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as XLA's Dot: https://www.tensorflow.org/xla/operation_semantics#dot

例:

%0 = mhlo.dot %arg0, %arg1 : (tensor<1x2xi32>, tensor<2x1xi32>) -> tensor<1x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dot_general (mhlo::DotGeneralOp)

DotGeneral operation

Computes dot products between slices of lhs and slices of rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dot_general

例:

%result = "mhlo.dot_general"(%lhs, %rhs) {
  dot_dimension_numbers = #mhlo.dot<
    lhs_batching_dimensions = [0],
    rhs_batching_dimensions = [0],
    lhs_contracting_dimensions = [2],
    rhs_contracting_dimensions = [1]
  >,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<2x2x2xi32>, tensor<2x2x2xi32>) -> tensor<2x2x2xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dot_dimension_numbers ::mlir::mhlo::DotDimensionNumbersAttr Attribute that models the dimension information for dot.
precision_config ::mlir::ArrayAttr Precision Config attribute
algorithm ::mlir::mhlo::DotAlgorithmAttr Attribute that models the algorithm constraints to use for computing dot.

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_broadcast_in_dim (mhlo::DynamicBroadcastInDimOp)

DynamicBroadcastInDim operation

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.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_broadcast_in_dim

例:

%operand = mhlo.constant dense<[[1, 2, 3]]> : tensor<1x3xi64>
%output_dimensions = mhlo.constant dense<[2, 3, 2]> : tensor<3xi64>
%result = "mhlo.dynamic_broadcast_in_dim"(%operand, %output_dimensions) {
  broadcast_dimensions = array<i64: 2, 1>,
  known_expanding_dimensions = array<i64: 0>,
  known_nonexpanding_dimensions = array<i64: 1>
} : (tensor<1x3xi64>, tensor<3xi64>) -> tensor<2x3x2xi64>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
broadcast_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
known_expanding_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
known_nonexpanding_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
output_dimensions 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_conv (mhlo::DynamicConvOp)

DynamicConv operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as ConvolutionOp except that padding is specified dynamically via d_padding : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convolution

例:

%result = "mhlo.dynamic_conv"(%lhs, %rhs, %d_padding) {
  window_strides = dense<4> : tensor<2xi64>,
  lhs_dilation = dense<2> : tensor<2xi64>,
  rhs_dilation = dense<1> : tensor<2xi64>,
  window_reversal = dense<false> : tensor<2xi1>,
  dimension_numbers = #mhlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi32>, tensor<3x3x1x1xi32>, tensor<2x2xi64>) -> tensor<1x2x2x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
lhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
rhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_reversal ::mlir::DenseElementsAttr constant boolean vector/tensor attribute
dimension_numbers ::mlir::mhlo::ConvDimensionNumbersAttr Structure of dimension information for conv op
feature_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
d_padding ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_gather (mhlo::DynamicGatherOp)

DynamicGather operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as GatherOp except that slice_sizes are specified dynamically: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#gather

例:

%result = "mhlo.dynamic_gather"(%operand, %start_indices, %slice_sizes) {
  dimension_numbers = #mhlo.gather<
    offset_dims = [2, 3],
    collapsed_slice_dims = [0],
    start_index_map = [0, 2],
    index_vector_dim = 2>,
  indices_are_sorted = false
} : (tensor<3x4x2xi32>, tensor<2x3x2xi64>, tensor<3xi64>) -> tensor<2x3x2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
slice_sizes statically shaped 1-dimensional integer tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_iota (mhlo::DynamicIotaOp)

DynamicIota operation

This operation is functionally identical to iota op, but the result shape is specified dynamically via output_shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_iota

例:

%0 = mhlo.dynamic_iota %arg0, dim = 0 : (tensor<1xindex>) -> tensor<4xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

オペランド説明
output_shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_pad (mhlo::DynamicPadOp)

DynamicPad operation

構文:

operation ::= `mhlo.dynamic_pad` operands attr-dict `:` functional-type(operands, results)

Dynamically Pads the operand , with amount of padding added at low-end/high-end/interior is passed through input tensors.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
edge_padding_low 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
edge_padding_high 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
interior_padding 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_reshape (mhlo::DynamicReshapeOp)

DynamicReshape operation

構文:

operation ::= `mhlo.dynamic_reshape` operands attr-dict `:` functional-type(operands, results)

This operation is functionally identical to reshape op, but the result shape is specified dynamically via output_shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_reshape

例:

%output_shape = mhlo.constant dense<[3, 2]> : tensor<2xi64>
%result = mhlo.dynamic_reshape %operand, %output_shape : (tensor<2x3xi64>, tensor<2xi64>) -> tensor<3x2xi64>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
output_shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_slice (mhlo::DynamicSliceOp)

DynamicSlice operation

Extracts a slice from the operand using dynamically-computed starting indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_slice

例:

%result = mhlo.dynamic_slice %operand, %start_indices0, %start_indices1, sizes = [2, 2]
  : (tensor<4x4xi32>, tensor<i64>, tensor<i64>) -> tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_update_slice (mhlo::DynamicUpdateSliceOp)

DynamicUpdateSlice operation

構文:

operation ::= `mhlo.dynamic_update_slice` operands attr-dict `:` functional-type(operands, results)

Produces a result tensor which is equal to the operand tensor except that the slice starting at start_indices is updated with the values in update .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_update_slice

例:

%result = mhlo.dynamic_update_slice %operand, %update, %start_indices0, %start_indices1
  : (tensor<4x4xi32>, tensor<2x2xi32>, tensor<i64>, tensor<i64>) -> tensor<4x4xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
update ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.einsum (mhlo::EinsumOp)

Einsum operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as TF's einsum: https://www.tensorflow.org/api_docs/python/tf/einsum

例:

%result = "mhlo.einsum"(%lhs, %rhs) {
  einsum_config = "ab,bc->ac"
} : (tensor<4x16xf32>, tensor<16x4xf32>) -> tensor<4x4xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
einsum_config ::mlir::StringAttr string attribute

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.erf (mhlo::ErfOp)

Erf operation

構文:

operation ::= `mhlo.erf` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise erf operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#erf

例:

%result = mhlo.erf %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.exponential (mhlo::ExpOp)

Exp operation

構文:

operation ::= `mhlo.exponential` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise exponential operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#exponential

例:

%result = mhlo.exponential %operand : tensor<2x2xf64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.exponential_minus_one (mhlo::Expm1Op)

Expm1 operation

構文:

operation ::= `mhlo.exponential_minus_one` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise exponential minus one operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#exponential_minus_one

例:

%result = mhlo.exponential_minus_one %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.fft (mhlo::FftOp)

Fft operation

Performs the forward and inverse Fourier transforms for real and complex inputs/outputs.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#fft

例:

%result = mhlo.fft %operand, type = FFT, length = [4] : (tensor<4xcomplex<f32>>) -> tensor<4xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
fft_type ::mlir::mhlo::FftTypeAttr XLA fast fourier transform type.
fft_length ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.floor (mhlo::FloorOp)

Floor operation

構文:

operation ::= `mhlo.floor` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise floor of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#floor

例:

%result = mhlo.floor %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.fusion (mhlo::FusionOp)

Fusion operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation consists of a group of basic ops (represented as a region attached to it). It serves as a hint to the backend that it is beneficial to emit the contained ops into a single loop nest or kernel.

Attributes:

属性MLIR Type説明
fusion_kind ::mlir::mhlo::FusionKindAttr fusion kind
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of Fusion

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

結果:

結果説明
results variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.gather (mhlo::GatherOp)

Gather operation

Gathers slices from operand tensor from offsets specified in start_indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#gather

例:

%result = "mhlo.gather"(%operand, %start_indices) {
  dimension_numbers = #stablehlo.gather<
    offset_dims = [3, 4],
    collapsed_slice_dims = [1],
    operand_batching_dims = [0],
    start_indices_batching_dims = [1],
    start_index_map = [2, 1],
    index_vector_dim = 3>,
  slice_sizes = dense<[0, 2, 2]> : tensor<3xi64>,
  indices_are_sorted = false
} : (tensor<2x3x4x2xi64>, tensor<2x2x3x2xi64>) -> tensor<2x2x3x2x2xi64>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.get_dimension_size (mhlo::GetDimensionSizeOp)

GetDimensionSize operation

Produces the size of the given dimension of the operand .

See https://github.com/openxla/stablehlo/blob/main/docs/spec.md#get_dimension_size

例:

%result = mhlo.get_dimension_size %operand, dim = 1 : (tensor<2x3xf32>) -> tensor<i32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» tensor of 32-bit signless integer values

mhlo.get_tuple_element (mhlo::GetTupleElementOp)

GetTupleElement operation

構文:

operation ::= `mhlo.get_tuple_element` $operand `[` $index `]` attr-dict `:` functional-type(operands, results)

Extracts element at index position of the operand tuple and produces a result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#get_tuple_element

例:

%result = mhlo.get_tuple_element %operand[0] : (tuple<tensor<2xf32>, tuple<tensor<i32>>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
index ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

オペランド説明
operand nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.if (mhlo::IfOp)

If operation

Produces the output from executing exactly one branch from true_branch or false_branch depending on the value of pred .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#if

Example: %result = "mhlo.if"(%pred) ({ "mhlo.return"(%result_true_branch) : (tensor ) -> () }, { "mhlo.return"(%result_false_branch) : (tensor ) -> () }) : (tensor ) -> tensor

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface

Operands:

オペランド説明
pred ranked tensor of pred (AKA boolean or 1-bit integer) values

結果:

結果説明
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.imag (mhlo::ImagOp)

Imag operation

構文:

operation ::= `mhlo.imag` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Extracts the imaginary part, element-wise, from the operand and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#imag

例:

%result = mhlo.imag %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.infeed (mhlo::InfeedOp)

Infeed operation

Reads data from the infeed and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#infeed

例:

%results:2 = "mhlo.infeed"(%token) {
  infeed_config = ""
} : (!mhlo.token) -> (tensor<3x3x3xi32>, !mhlo.token)

Attributes:

属性MLIR Type説明
infeed_config ::mlir::StringAttr string attribute
layout ::mlir::ArrayAttr array attribute

Operands:

オペランド説明
tokenトークン

結果:

結果説明
«unnamed» variadic of statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.iota (mhlo::IotaOp)

Iota operation

Fills an output tensor with values in increasing order starting from zero along the iota_dimension dimension.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#iota

例:

%output = mhlo.iota dim = 0 : tensor<4x5xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

結果:

結果説明
output statically shaped tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.is_finite (mhlo::IsFiniteOp)

IsFinite operation

構文:

operation ::= `mhlo.is_finite` $x attr-dict `:` functional-type(operands, results)

Performs element-wise check whether the value in x is finite (ie is neither +Inf, -Inf, nor NaN) and produces a y tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#is_finite

例:

%y = mhlo.is_finite %x : (tensor<7xf32>) -> tensor<7xi1>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
x ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

結果:

結果説明
y ranked tensor of pred (AKA boolean or 1-bit integer) values

mhlo.log (mhlo::LogOp)

Log operation

構文:

operation ::= `mhlo.log` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logarithm operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#log

例:

%result = mhlo.log %operand : tensor<2x2xf64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.log_plus_one (mhlo::Log1pOp)

Log1p operation

構文:

operation ::= `mhlo.log_plus_one` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logarithm plus one operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#log_plus_one

例:

%result = mhlo.log_plus_one %operand : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.logistic (mhlo::LogisticOp)

Logistic operation

構文:

operation ::= `mhlo.logistic` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logistic operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#logistic

例:

%result = mhlo.logistic %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.map (mhlo::MapOp)

Map operation

Applies a map function computation to inputs along the dimensions and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#map

例:

%result = "mhlo.map"(%input0, %input1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.multiply %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  dimensions = dense<[0, 1]> : tensor<2xi64>
} : (tensor<2x2xi32>, tensor<2x2xi32>) -> tensor<2x2xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameOperandsAndResultShape , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

属性MLIR Type説明
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.maximum (mhlo::MaxOp)

Max operation

構文:

operation ::= `mhlo.maximum` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise max operation on tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#maximum

例:

%result = mhlo.maximum %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.minimum (mhlo::MinOp)

Min operation

構文:

operation ::= `mhlo.minimum` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise min operation on tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#minimum

例:

%result = mhlo.minimum %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.minimum_broadcast_shapes (mhlo::MinimumBroadcastShapesOp)

Minimizes the rank of two or more shapes to be broadcasted

構文:

operation ::= `mhlo.minimum_broadcast_shapes` $shapes attr-dict `:` type($shapes) `->` type($results)

Given two or more 1D tensors representing shapes, returns one 1D tensor for each operand, where operand i corresponds to output i .

The returned tensors have the property that they specify a shape which is a reshape of the corresponding input shape, and the broadcasted output shape (using shape::BroadcastOp) of the returned shapes is a reshape of the broadcasted output shape of the input shapes. Among all possibilities with this property, the one is chosen which minimizes the rank of each returned shape.

The general idea of this op is that it can be used for ops which have a broadcasting semantic to operate on shapes with a possibly smaller rank while preserving equivalence of the computed values. After computing the result of the op using reshaped operands, the result can be reshaped to the result that would have been originally computed.

Here is an example with two input shapes:

mhlo.minimum_broadcast_shapes [1, 2, 3, 1, 2, 1],
                                 [1, 1, 1, 2, 3] -> [6, 2, 1], [2, 3]

The broadcasted output shape of the operands is [1, 2, 3, 1, 2, 3], the broadcasted output shape of the outputs is [6, 2, 3]. These two shapes are reshapes of each other, and also each output is a reshape of the corresponding input.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
shapes variadic of 1D tensor of index values

結果:

結果説明
results variadic of 1D tensor of index values

mhlo.multiply (mhlo::MulOp)

Mul operation

構文:

operation ::= `mhlo.multiply` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise product of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#multiply

例:

%result = mhlo.multiply %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.negate (mhlo::NegOp)

Neg operation

構文:

operation ::= `mhlo.negate` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise negation of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#negate

例:

%result = mhlo.negate %operand : tensor<2x3xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.not (mhlo::NotOp)

Not operation

構文:

operation ::= `mhlo.not` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise NOT of tensor operand of type integer and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#not

例:

%result = mhlo.not %operand : tensor<5x3x1xi1>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.optimization_barrier (mhlo::OptimizationBarrierOp)

OptimizationBarrier operation

構文:

operation ::= `mhlo.optimization_barrier` attr-dict ($operand^ `:` custom<PairwiseOpType>(type($operand), type($result))):(`(` `)`)?

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, ie result = operand .

See https://github.com/openxla/stablehlo/blob/main/docs/spec.md#optimization_barrier

例:

%result0, %result1 = mhlo.optimization_barrier %operand0, %operand1 : tensor<f32>, tensor<f32>

Traits: AlwaysSpeculatableImplTrait , HLO_PairwiseSameOperandAndResultType

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

結果:

結果説明
result variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.or (mhlo::OrOp)

Or operation

構文:

operation ::= `mhlo.or` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise OR of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#or

例:

%result = mhlo.or %lhs, %rhs : tensor<2xi1>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.outfeed (mhlo::OutfeedOp)

Outfeed operation

Writes inputs to the outfeed and produces a result token.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#outfeed

例:

%result = "mhlo.outfeed"(%input0, %token) {
  outfeed_config = ""
} : (tensor<3x3x3xi32>, !mhlo.token) -> !mhlo.token

Interfaces: InferTypeOpInterface

Attributes:

属性MLIR Type説明
outfeed_config ::mlir::StringAttr string attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
tokenトークン

結果:

結果説明
«unnamed»トークン

mhlo.pad (mhlo::PadOp)

Pad operation

Expands operand by padding around the tensor as well as between the elements of the tensor with the given padding_value .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#pad

例:

%0 = mhlo.pad %arg0, %arg1, low = [0, 1], high = [2, 1], interior = [1, 2]
  : (tensor<2x3xi32>, tensor<i32>) -> tensor<5x9xi32>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
edge_padding_low ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
edge_padding_high ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
interior_padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.partition_id (mhlo::PartitionIdOp)

PartitionId operation

構文:

operation ::= `mhlo.partition_id` attr-dict `:` type(results)

Produces partition_id of the current process.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#partition_id

例:

%result = mhlo.partition_id : tensor<ui32>

Interfaces: InferTypeOpInterface

結果:

結果説明
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.popcnt (mhlo::PopulationCountOp)

PopulationCount operation

構文:

operation ::= `mhlo.popcnt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise count of the number of bits set in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#popcnt

例:

%result = mhlo.popcnt %operand : tensor<4xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.power (mhlo::PowOp)

Pow operation

構文:

operation ::= `mhlo.power` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise exponentiation of lhs tensor by rhs tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#power

例:

%result = mhlo.power %lhs, %rhs : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.real (mhlo::RealOp)

Real operation

構文:

operation ::= `mhlo.real` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Extracts the real part, element-wise, from the operand and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#real

例:

%result = mhlo.real %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.real_dynamic_slice (mhlo::RealDynamicSliceOp)

RealDynamicSlice operation

構文:

operation ::= `mhlo.real_dynamic_slice` operands attr-dict `:` functional-type(operands, results)

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as SliceOp except that start_indices , limit_indices and strides are specified dynamically: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#slice

例:

%result = mhlo.real_dynamic_slice %operand,
            %start_indices, %limit_indices, %strides
       : (tensor<256x?xf32>, tensor<2xindex>, tensor<2xindex>, tensor<2xindex>) -> tensor<256x?xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
limit_indices 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
strides 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.recv (mhlo::RecvOp)

Recv operation

Receives data from a channel with channel_id and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#recv

例:

%results:2 = "mhlo.recv"(%token) {
  // channel_id = 5 : i64,
  // channel_type = #stablehlo<channel_type HOST_TO_DEVICE>,
  channel_handle = #mhlo.channel_handle<handle = 5, type = 3>,
  is_host_transfer = true
} : (!mhlo.token) -> (tensor<3x4xi32>, !mhlo.token)

Attributes:

属性MLIR Type説明
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
tokenトークン

結果:

結果説明
«unnamed» variadic of statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.reduce (mhlo::ReduceOp)

Reduce operation

Applies a reduction function body to inputs and init_values along the dimensions and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce

例:

%result = "mhlo.reduce"(%input, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0) : (tensor<i32>) -> ()
}) {
  dimensions = dense<1> : tensor<1xi64>
} : (tensor<1x6xi32>, tensor<i32>) -> tensor<1xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameVariadicOperandSize , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

属性MLIR Type説明
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_values variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.reduce_precision (mhlo::ReducePrecisionOp)

ReducePrecision operation

構文:

operation ::= `mhlo.reduce_precision` $operand `,` `format` `=` custom<ExponentMantissa>($exponent_bits, $mantissa_bits)
              attr-dict `:` custom<SameOperandsAndResultType>(type($operand), type($output))

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.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_precision

例:

%output = mhlo.reduce_precision %operand, format = e5m2 : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
exponent_bits ::mlir::IntegerAttr 32-bit signless integer attribute whose value is positive
mantissa_bits ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

結果:

結果説明
output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.reduce_scatter (mhlo::ReduceScatterOp)

ReduceScatter operation

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 .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_scatter

例:

%result = "mhlo.reduce_scatter"(%operand) ({
  ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>):
  %0 = mhlo.add %arg0, %arg1 : tensor<f32>
  mhlo.return %0 : tensor<f32>
}) {
  scatter_dimension = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<2x4xf32>) -> tensor<2x2xf32>

Attributes:

属性MLIR Type説明
scatter_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
use_global_device_ids ::mlir::UnitAttr unit attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.reduce_window (mhlo::ReduceWindowOp)

ReduceWindow operation

Applies a reduction function body to windows of inputs and init_values and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_window

例:

%result = "mhlo.reduce_window"(%input, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.add %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  window_dimensions = dense<[2, 1]> : tensor<2xi64>,
  window_strides = dense<[4, 1]> : tensor<2xi64>,
  base_dilations = dense<[2, 1]> : tensor<2xi64>,
  window_dilations = dense<[3, 1]> : tensor<2xi64>,
  padding = dense<[[2, 1], [0, 0]]> : tensor<2x2xi64>
} : (tensor<3x2xi32>, tensor<i32>) -> tensor<2x2xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameVariadicOperandSize , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

属性MLIR Type説明
window_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
base_dilations ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_dilations ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_values variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.remainder (mhlo::RemOp)

Rem operation

構文:

operation ::= `mhlo.remainder` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise remainder of dividend lhs and divisor rhs tensors and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#remainder

例:

%result = mhlo.remainder %lhs, %rhs : tensor<4xi64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.replica_id (mhlo::ReplicaIdOp)

ReplicaId operation

構文:

operation ::= `mhlo.replica_id` attr-dict `:` type(results)

Produces replica_id of the current process.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#replica_id

例:

%result = mhlo.replica_id : tensor<ui32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

結果:

結果説明
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.reshape (mhlo::ReshapeOp)

Reshape operation

構文:

operation ::= `mhlo.reshape` operands attr-dict `:` functional-type(operands, results)

Performs reshape of operand tensor to a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reshape

例:

%result = mhlo.reshape %operand : (tensor<2xf32>) -> tensor<1x2xf32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.return (mhlo::ReturnOp)

_This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/425

Informally, this operation serves as a terminator for regions defined by
the StableHLO ops. Non-StableHLO ops, e.g. `func.func`, have their own
terminators, e.g. `func.return`.

Example:

    ```mlir
    %result = "mhlo.reduce"(%input, %init_value) ({
      ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
        %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
        "mhlo.return"(%0) : (tensor<i32>) -> ()
    }) {
      dimensions = dense<1> : tensor<1xi64>
    } : (tensor<1x6xi32>, tensor<i32>) -> tensor<1xi32>
    ```_


Syntax:

```

operation ::= mhlo.return $results attr-dict ( : type($results)^)?



Traits: `AlwaysSpeculatableImplTrait`, `Terminator`

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

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `results` | variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values


### `mhlo.reverse` (mhlo::ReverseOp)

_Reverse operation_

Reverses the order of elements in the `operand` along the specified
`dimensions` and produces a `result` tensor.

See:
<a href="https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reverse">https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reverse</a>

Example:
```mlir
%result = mhlo.reverse %operand, dims = [1] : tensor<3x2xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.rng (mhlo::RngOp)

Rng operation

Generates random numbers using the rng_distribution algorithm and produces a result tensor of a given shape shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rng

例:

%result = mhlo.rng %a, %b, %shape, distribution = NORMAL : (tensor<i32>, tensor<i32>, tensor<2xi64>) -> tensor<3x3xi32>

Traits: InferTensorType

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

属性MLIR Type説明
rng_distribution ::mlir::mhlo::RngDistributionAttr XLA PRNG distribution to be used.

Operands:

オペランド説明
a 0D tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
b 0D tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.rng_bit_generator (mhlo::RngBitGeneratorOp)

RngBitGenerator operation

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 .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rng_bit_generator

例:

%output_state, %output = mhlo.rng_bit_generator %initial_state, algorithm = THREE_FRY : (tensor<2xui64>) -> (tensor<2xui64>, tensor<2x2xui64>)

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
rng_algorithm ::mlir::mhlo::RngAlgorithmAttr XLA PRNG algorithm to be used.

Operands:

オペランド説明
initial_state ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

結果:

結果説明
output_state ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
output statically shaped tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.round_nearest_afz (mhlo::RoundOp)

Round operation

構文:

operation ::= `mhlo.round_nearest_afz` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise rounding towards the nearest integer, breaking ties away from zero, on the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#round_nearest_afz

例:

%result = mhlo.round_nearest_afz %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.round_nearest_even (mhlo::RoundNearestEvenOp)

RoundNearestEven operation

構文:

operation ::= `mhlo.round_nearest_even` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise rounding towards the nearest integer, breaking ties towards the even integer, on the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#round_nearest_even

例:

%result = mhlo.round_nearest_even %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.rsqrt (mhlo::RsqrtOp)

Rsqrt operation

構文:

operation ::= `mhlo.rsqrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise reciprocal square root operation on operand tensor and produces a result tensor, implementing the rSqrt operation from the IEEE-754 specification.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rsqrt

例:

%result = mhlo.rsqrt %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.scatter (mhlo::ScatterOp)

Scatter operation

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 .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#scatter

例:

%result = "mhlo.scatter"(%input, %scatter_indices, %update) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.add %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  scatter_dimension_numbers = #mhlo.scatter<
    update_window_dims = [3, 4],
    inserted_window_dims = [1],
    input_batching_dims = [0],
    scatter_indices_batching_dims = [1],
    scatter_dims_to_operand_dims = [2, 1],
    index_vector_dim = 3>,
  indices_are_sorted = false,
  unique_indices = false
} : (tensor<2x3x4x2xi64>, tensor<2x2x3x2xi64>, tensor<2x2x3x2x2xi64>) -> tensor<2x3x4x2xi64>

Traits: RecursiveMemoryEffects , SameVariadicOperandSize

Interfaces: InferTypeOpInterface

Attributes:

属性MLIR Type説明
scatter_dimension_numbers ::mlir::mhlo::ScatterDimensionNumbersAttr Attribute that models the dimension information for scatter
indices_are_sorted ::mlir::BoolAttr bool attribute
unique_indices ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
scatter_indices ranked tensor of integer or index values
updates variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.select (mhlo::SelectOp)

Select operation

構文:

operation ::= `mhlo.select` operands attr-dict `:`
              custom<SelectOpType>(type($pred), type($on_true), type($on_false), type($result))

Produces a result tensor where each element is selected from on_true or on_false tensor based on the value of the corresponding element of pred .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#select

例:

%result = mhlo.select %pred, %on_true, %on_false : tensor<2x2xi1>, tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_BroadcastingElementwise , InferTensorType

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
pred ranked tensor of pred (AKA boolean or 1-bit integer) values
on_true ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
on_false ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.select_and_scatter (mhlo::SelectAndScatterOp)

SelectAndScatter operation

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.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#select_and_scatter

例:

%result = "mhlo.select_and_scatter"(%operand, %source, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GE>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0) : (tensor<i32>) -> ()
}) {
  window_dimensions = dense<[3, 1]> : tensor<2xi64>,
  window_strides = dense<[2, 1]> : tensor<2xi64>,
  padding = dense<[[0, 1], [0, 0]]> : tensor<2x2xi64>
} : (tensor<4x2xi32>, tensor<2x2xi32>, tensor<i32>) -> tensor<4x2xi32>

Traits: RecursiveMemoryEffects

Interfaces: InferTypeOpInterface

Attributes:

属性MLIR Type説明
window_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
source ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.send (mhlo::SendOp)

Send operation

Sends inputs to a channel channel_id and produces a result token.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#send

例:

%result = "mhlo.send"(%operand, %token) {
  // channel_id = 5 : i64,
  // channel_type = #stablehlo<channel_type DEVICE_TO_HOST>,
  channel_handle = #mhlo.channel_handle<handle = 5, type = 2>,
  is_host_transfer = true
} : (tensor<3x4xi32>, !mhlo.token) -> !mhlo.token

Interfaces: InferTypeOpInterface

Attributes:

属性MLIR Type説明
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
tokenトークン

結果:

結果説明
«unnamed»トークン

mhlo.set_dimension_size (mhlo::SetDimensionSizeOp)

SetDimensionSize operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as XLA's SetDimensionSize: https://www.tensorflow.org/xla/operation_semantics#setdimensionsize

例:

%0 = mhlo.set_dimension_size %arg0, %arg1, dim = 1 : (tensor<4x2xf32>, tensor<i32>) -> tensor<4x2xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
size tensor of 32-bit signless integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.shift_left (mhlo::ShiftLeftOp)

ShiftLeft operation

構文:

operation ::= `mhlo.shift_left` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise left-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_left

例:

%result = mhlo.shift_left %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.shift_right_arithmetic (mhlo::ShiftRightArithmeticOp)

ShiftRightArithmetic operation

構文:

operation ::= `mhlo.shift_right_arithmetic` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise arithmetic right-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_right_arithmetic

例:

%result = mhlo.shift_right_arithmetic %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.shift_right_logical (mhlo::ShiftRightLogicalOp)

ShiftRightLogical operation

構文:

operation ::= `mhlo.shift_right_logical` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise logical right-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_right_logical

例:

%result = mhlo.shift_right_logical %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.sign (mhlo::SignOp)

Sign operation

構文:

operation ::= `mhlo.sign` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Returns the sign of the operand element-wise and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sign

例:

%result = mhlo.sign %operand : tensor<7xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.sine (mhlo::SineOp)

Sine operation

構文:

operation ::= `mhlo.sine` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise sine operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sine

例:

%result = mhlo.sine %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.slice (mhlo::SliceOp)

Slice operation

Extracts a slice from the operand using statically-computed starting indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#slice

例:

%result = "mhlo.slice" (%operand) {
  start_indices = dense<[1, 2]> : tensor<2xi64>,
  limit_indices = dense<[3, 4]> : tensor<2xi64>,
  strides = dense<1> : tensor<2xi64>
} : (tensor<3x4xi64>) -> tensor<2x2xi64>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
start_indices ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
limit_indices ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sort (mhlo::SortOp)

Sort operation

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 .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sort

例:

%result0, %result1 = "mhlo.sort"(%input0, %input1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>, %arg2: tensor<i32>, %arg3: tensor<i32>):
    %predicate = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GT>
      } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%predicate) : (tensor<i1>) -> ()
}) {
  dimension = 0 : i64,
  is_stable = true
} : (tensor<2x3xi32>, tensor<2x3xi32>) -> (tensor<2x3xi32>, tensor<2x3xi32>)

Traits: InferTensorType , RecursiveMemoryEffects , SameOperandsAndResultShape

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

属性MLIR Type説明
dimension ::mlir::IntegerAttr 64-bit signless integer attribute
is_stable ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sparse_dot (mhlo::SparseDotOp)

Sparse dot operation

Similar to dot_general operation, with one or both of the operands being sparse. An additional argument provides sparsity meta information. Disclaimer: this op is experimental / a work in progress.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
lhs_sparsity ::mlir::mhlo::SparsityDescriptorAttr Describes structured (N:M) sparsity configuration
rhs_sparsity ::mlir::mhlo::SparsityDescriptorAttr Describes structured (N:M) sparsity configuration
dot_dimension_numbers ::mlir::mhlo::DotDimensionNumbersAttr Attribute that models the dimension information for dot.
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

オペランド説明
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
meta variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sqrt (mhlo::SqrtOp)

Sqrt operation

構文:

operation ::= `mhlo.sqrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise square root operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sqrt

例:

%result = mhlo.sqrt %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.stochastic_convert (mhlo::StochasticConvertOp)

StochasticConvert operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/295

Informally, this operation performs element-wise conversion of values from a bigger type to a smaller one with stochastic rounding using the random number passed in.

Traits: AlwaysSpeculatableImplTrait , Elementwise

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
random ranked tensor of 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.subtract (mhlo::SubtractOp)

Subtract operation

構文:

operation ::= `mhlo.subtract` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise subtraction of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#subtract

例:

%result = mhlo.subtract %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.tan (mhlo::TanOp)

Tan operation

構文:

operation ::= `mhlo.tan` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/954

Informally, this operation returns Tan(operand) element-wise.

例:

%0 = mhlo.tan %arg0 : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.tanh (mhlo::TanhOp)

Tanh operation

構文:

operation ::= `mhlo.tanh` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise hyperbolic tangent operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#tanh

例:

%result = mhlo.tanh %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.topk (mhlo::TopKOp)

TopK operation

構文:

operation ::= `mhlo.topk` `(`$operand `,` `k` `=` $k (`,` `largest` `=` $largest^)? `)` attr-dict `:`
              type($operand) `->` `(`type($values)`,` type($indices)`)`

Returns top k values and their indices, along the last dimension of the operand if largest=true or the bottom k values if largest=false .

See: https://www.tensorflow.org/xla/operation_semantics#top-k

例:

%values, %indices = mhlo.topk(%operand, k=5, largest=true)
  : tensor<100xf32> -> (tensor<5xf32>, tensor<5xi32>)

Traits: InferTensorType , RecursiveMemoryEffects

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

属性MLIR Type説明
k ::mlir::IntegerAttr 64-bit signless integer attribute
largest ::mlir::BoolAttr bool attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
values ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
indices ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.torch_index_select (mhlo::TorchIndexSelectOp)

TorchIndexSelect operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as PyTorch's index_select, augmented with support for batch dimensions: https://pytorch.org/docs/stable/generated/torch.index_select.html

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.

例:

%result = "mhlo.torch_index_select"(%operand, %index) {
  dim = 2 : i64,
  batch_dims = 1 : i64
} : (tensor<8x128x3072x64xf32>, tensor<8x16x1024xi32>) -> tensor<8x128x16x1024x64xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
dim ::mlir::IntegerAttr 64-bit signless integer attribute
batch_dims ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
index ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.trace (mhlo::TraceOp)

Trace operation

構文:

operation ::= `mhlo.trace` $operand `,` $tag attr-dict `:` type($operand)

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/604

It is not used by JAX, PyTorch or TensorFlow, so it looks like we should've classified it as "Private to XLA" and not included it in StableHLO in the first place. With that in mind, its semantics will not be documented here.

例:

mhlo.trace %arg0, "In test code." : tensor<5x1x5xi32>

Attributes:

属性MLIR Type説明
tag ::mlir::StringAttr string attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.transpose (mhlo::TransposeOp)

Transpose operation

Permutes the dimensions of operand tensor using permutation and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#transpose

例:

%0 = mhlo.transpose %arg0, dims = [2, 1, 0] : (tensor<1x2x3xi32>) -> tensor<3x2x1xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
permutation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.triangular_solve (mhlo::TriangularSolveOp)

TriangularSolve operation

Solves batches of systems of linear equations with lower or upper triangular coefficient matrices.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#triangular_solve

例:

%result = "mhlo.triangular_solve"(%a, %b) {
  left_side = true,
  lower = true,
  unit_diagonal = false,
  transpose_a = #stablehlo<transpose NO_TRANSPOSE>
} : (tensor<3x3xf32>, tensor<3x3xf32>) -> tensor<3x3xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

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

Effects: MemoryEffects::Effect{}

Attributes:

属性MLIR Type説明
left_side ::mlir::BoolAttr bool attribute
lower ::mlir::BoolAttr bool attribute
unit_diagonal ::mlir::BoolAttr bool attribute
transpose_a ::mlir::mhlo::TransposeAttr Transpose options

Operands:

オペランド説明
a ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values
b ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

結果:

結果説明
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.tuple (mhlo::TupleOp)

Tuple operation

構文:

operation ::= `mhlo.tuple` $val attr-dict `:` custom<TupleOpType>(type($val), type($result))

Produces a result tuple from values val .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#tuple

例:

%result = mhlo.tuple %val0, %val1 : tuple<tensor<2xf32>, tuple<tensor<i32>>>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
val variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

結果:

結果説明
result nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.uniform_dequantize (mhlo::UniformDequantizeOp)

UniformDequantize operation

構文:

operation ::= `mhlo.uniform_dequantize` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise conversion of quantized tensor operand to a floating-point tensor result according to the quantization parameters defined by the operand type.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#uniform_dequantize

例:

%result = mhlo.uniform_dequantize %operand : (tensor<16x16x!quant.uniform<i8:f32, 34.0:16>>) -> tensor<16x16xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , InferTensorType , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.uniform_quantize (mhlo::UniformQuantizeOp)

UniformQuantize operation

構文:

operation ::= `mhlo.uniform_quantize` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

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.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#uniform_quantize

例:

%result = mhlo.uniform_quantize %operand : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

結果:

結果説明
result ranked tensor of 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.while (mhlo::WhileOp)

While operation

Produces the output from executing body function 0 or more times while the cond function outputs true .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#while

例:

%results0, %results1 = "mhlo.while"(%operand0, %operand1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction LT>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %constant0) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0, %arg1) : (tensor<i32>, tensor<i32>) -> ()
}) : (tensor<i32>, tensor<i32>) -> (tensor<i32>, tensor<i32>)

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface , OpAsmOpInterface

Operands:

オペランド説明
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

結果:

結果説明
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.xla.rng_get_and_update_state (mhlo::XlaRngGetAndUpdateStateOp)

XlaRngGetAndUpdateState operation

構文:

operation ::= `mhlo.xla.rng_get_and_update_state` attr-dict

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation represents the change of the global random number generator state for rng instructions. The global state is incremented by delta and the old state is returned.

The output is currently defined for a single output type. If this changes in the future to support multiple types, lowering to use of a global memref must ensure that a single memref is still used and updated appropriately.

Interfaces: InferTypeOpInterface

Attributes:

属性MLIR Type説明
delta ::mlir::IntegerAttr 64-bit signless integer attribute

結果:

結果説明
«unnamed» statically shaped tensor of 64-bit unsigned integer values

mhlo.xor (mhlo::XorOp)

Xor operation

構文:

operation ::= `mhlo.xor` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise XOR of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#xor

例:

%result = mhlo.xor %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

オペランド説明
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

結果:

結果説明
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Attributes

ArgResultAliasAttr

Attribute that models the alias relationship of entry function argument

This attribute captures the alias relationship of an MHLO main function argument to one of the results, denoted by resultIndex . The argTupleIndices and resultTupleIndices are used to index into nested tuples in operand and result respectively. If isMustAlias is true then the operand-result pair must alias.

This is meant to be used as an attribute on a function argument in MHLO. For example, in the following code it expresses that %arg1 may alias 0-th result.

func @main(%arg0: tensor<2xf32>, %arg1: tensor<3xf32> {mhlo.result_alias =
    mhlo.result_alias<result_index = [2], ...>}
  ) -> tensor<2xf32>, tensor<3xf32> {
  // function body ...
}

Parameters:

Parameter C++ type説明
argTupleIndices ::llvm::ArrayRef<int64_t>寸法
resultIndex int64_t
resultTupleIndices ::llvm::ArrayRef<int64_t>寸法
isMustAlias bool

ChannelHandleAttr

two 64-bit integers 'handle' and 'type'

構文:

#mhlo.channel_handle<
  int64_t,   # handle
  int64_t   # type
>

Parameters:

Parameter C++ type説明
ハンドルint64_t
タイプint64_t

ComparisonDirectionAttr

Which comparison operation to perform.

構文:

#mhlo.comparison_direction<
  ::mlir::mhlo::ComparisonDirection   # value
>

Enum cases:

  • EQ ( EQ )
  • NE ( NE )
  • GE ( GE )
  • GT ( GT )
  • LE ( LE )
  • LT ( LT )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::ComparisonDirection an enum of type ComparisonDirection

ComparisonTypeAttr

Which comparison type to use.

構文:

#mhlo.comparison_type<
  ::mlir::mhlo::ComparisonType   # value
>

Enum cases:

  • NOTYPE ( NOTYPE )
  • FLOAT ( FLOAT )
  • TOTALORDER ( TOTALORDER )
  • SIGNED ( SIGNED )
  • UNSIGNED ( UNSIGNED )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::ComparisonType an enum of type ComparisonType

ConvDimensionNumbersAttr

Structure of dimension information for conv op

Parameters:

Parameter C++ type説明
inputBatchDimension int64_t
inputFeatureDimension int64_t
inputSpatialDimensions ::llvm::ArrayRef<int64_t>寸法
kernelInputFeatureDimension int64_t
kernelOutputFeatureDimension int64_t
kernelSpatialDimensions ::llvm::ArrayRef<int64_t>寸法
outputBatchDimension int64_t
outputFeatureDimension int64_t
outputSpatialDimensions ::llvm::ArrayRef<int64_t>寸法

CrossProgramPrefetchAttr

Argument that is prefetched from another program

構文:

#mhlo.cross_program_prefetch<
  int64_t,   # parameter
  ::llvm::ArrayRef<int64_t>,   # indices
  std::optional<int64_t>   # offset
>

This attribute captures an argument that is prefetched from another program. For a given CrossProgramPrefetchAttr , parameter tells us which argument of the main function of the module is prefetched, and indices is a shape index telling us what subshape of that argument is prefetched.

A shape has a subshape iff it is a tuple. In that case, the subshape of the tuple by indices is the shape achieved after indexing by each element of indices in turn. For example, the [1,0] subshape of tuple<tuple<token, token>, tuple<tensor<i32>, token>> is tensor<i32> .

An empty value for indices means the whole shape is prefetched.

例えば、

module attributes { mhlo.cross_program_prefetch = [ #mhlo.cross_program_prefetch< parameter = 0, indices = [0]> ]} {
  func.func @copy(%arg0 : tuple<tensor<2x3xi32>, tensor<i32>>) -> tuple<tensor<2x3xi32>, tensor<i32>> {
    %0 = "mhlo.copy"(%arg0) {is_cross_program_prefetch}
    return %0 : tuple<tensor<2x3xi32>, tensor<i32>>
  }
  func.func @main(%arg0 : tuple<tensor<2x3xi32>, tensor<i32>>) -> tuple<tensor<2x3xi32>, tensor<i32>> {
    %1 = "mhlo.async_start"(%arg0) {called_computation=@copy}
    %2 = "mhlo.async_done"(%1) {called_computation=@copy}
    return %2 : tuple<tensor<2x3xi32>, tensor<i32>>
  }
}

The parameter = 0 tells us that the async copy of the 0 th parameter is a cross_program_prefetch , while the index of [0] tells us that the 0 th element of the tuple is prefetched while the other element of the tuple is not.

Parameters:

Parameter C++ type説明
パラメータint64_t
インデックス::llvm::ArrayRef<int64_t>寸法
オフセットstd::optional<int64_t>

CustomCallScheduleAttr

Specifies the desired schedule for the custom-call.

構文:

#mhlo.custom_call_schedule<
  ::mlir::mhlo::CustomCallSchedule   # value
>

Enum cases:

  • NONE ( NONE )
  • LATEST ( LATEST )
  • EARLIEST ( EARLIEST )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::CustomCallSchedule an enum of type CustomCallSchedule

DequantizeModeAttr

Dequantization mode. Only MIN_COMBINED is supported.

構文:

#mhlo.dequantize_mode<
  ::mlir::mhlo::DequantizeMode   # value
>

Enum cases:

  • MIN_COMBINED ( MIN_COMBINED )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::DequantizeMode an enum of type DequantizeMode

DomainKindAttr

Kind of domain metatdata attached to an HLO domain.

構文:

#mhlo.kind<
  ::mlir::mhlo::DomainKind   # value
>

Enum cases:

  • sharding ( sharding )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::DomainKind an enum of type DomainKind

DotAlgorithmAttr

Attribute that models the algorithm constraints to use for computing dot.

構文:

#mhlo.dot_algorithm<
  Type,   # lhsPrecisionType
  Type,   # rhsPrecisionType
  Type,   # accumulationType
  int64_t,   # lhsComponentCount
  int64_t,   # rhsComponentCount
  int64_t,   # numPrimitiveOperations
  bool   # allowImpreciseAccumulation
>

Parameters:

Parameter C++ type説明
lhsPrecisionType Type
rhsPrecisionType Type
accumulationType Type
lhsComponentCount int64_t
rhsComponentCount int64_t
numPrimitiveOperations int64_t
allowImpreciseAccumulation bool

DotDimensionNumbersAttr

Attribute that models the dimension information for dot.

Parameters:

Parameter C++ type説明
lhsBatchingDimensions ::llvm::ArrayRef<int64_t>寸法
rhsBatchingDimensions ::llvm::ArrayRef<int64_t>寸法
lhsContractingDimensions ::llvm::ArrayRef<int64_t>寸法
rhsContractingDimensions ::llvm::ArrayRef<int64_t>寸法

FftTypeAttr

XLA fast fourier transform type.

構文:

#mhlo.fft_type<
  ::mlir::mhlo::FftType   # value
>

Enum cases:

  • FFT ( FFT )
  • IFFT ( IFFT )
  • RFFT ( RFFT )
  • IRFFT ( IRFFT )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::FftType an enum of type FftType

FusionKindAttr

fusion kind

構文:

#mhlo.fusion_kind<
  ::mlir::mhlo::FusionKind   # value
>

Enum cases:

  • kLoop ( kLoop )
  • kInput ( kInput )
  • kOutput ( kOutput )
  • kCustom ( kCustom )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::FusionKind an enum of type FusionKind

GatherDimensionNumbersAttr

Attribute that models the dimension information for gather

Parameters:

Parameter C++ type説明
offsetDims ::llvm::ArrayRef<int64_t>寸法
collapsedSliceDims ::llvm::ArrayRef<int64_t>寸法
operandBatchingDims ::llvm::ArrayRef<int64_t>寸法
startIndicesBatchingDims ::llvm::ArrayRef<int64_t>寸法
startIndexMap ::llvm::ArrayRef<int64_t>寸法
indexVectorDim int64_t

OutputOperandAliasAttr

Attribute that models the alias relationship of output and operand of a CustomCall op

構文:

#mhlo.output_operand_alias<
  ::llvm::ArrayRef<int64_t>,   # outputTupleIndices
  int64_t,   # operandIndex
  ::llvm::ArrayRef<int64_t>   # operandTupleIndices
>

This attribute captures the alias relationship of the output to one of the operands for a CustomCall op, denoted by operand_index . The output_tuple_indices and operand_tuple_indices are used to index into output and operand types. These indices lists are empty if the corresponding types are not tuple types, and can be arbitrarily long in case of arbitrarily nested tuple types.

See https://www.tensorflow.org/xla/aliasing

Example when used as array with in mhlo.custom-call:

%0 = "mhlo.custom_call"(%arg0, %arg1) {
  // other attributes
  output_operand_alias = [
    #mhlo.output_operand_alias<output_tuple_indices = [0],
                               operand_index = 0,
                               operand_tuple_indices = [1]>
  ]
} : (tuple<tensor<1x1xf32>, tensor<2x3xf32>>, tensor<5x5xf32>) -> tuple<tensor<2x3xf32>>

The output and the 0th operand are both tuples. The aliasing shows the
relationship between the 0th element in output tuple with the 1st element in
the 0th operand. And both of them are of the same type: tensor<2x3xf32>.

Parameters:

Parameter C++ type説明
outputTupleIndices ::llvm::ArrayRef<int64_t>寸法
operandIndex int64_t
operandTupleIndices ::llvm::ArrayRef<int64_t>寸法

PrecisionAttr

XLA precision for an operand. Has backend specific meaning.

構文:

#mhlo.precision<
  ::mlir::mhlo::Precision   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • HIGH ( HIGH )
  • HIGHEST ( HIGHEST )
  • PACKED_NIBBLE ( PACKED_NIBBLE )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::Precision an enum of type Precision

RngAlgorithmAttr

XLA PRNG algorithm to be used.

構文:

#mhlo.rng_algorithm<
  ::mlir::mhlo::RngAlgorithm   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • THREE_FRY ( THREE_FRY )
  • PHILOX ( PHILOX )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::RngAlgorithm an enum of type RngAlgorithm

RngDistributionAttr

XLA PRNG distribution to be used.

構文:

#mhlo.rng_distribution<
  ::mlir::mhlo::RngDistribution   # value
>

Enum cases:

  • UNIFORM ( UNIFORM )
  • NORMAL ( NORMAL )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::RngDistribution an enum of type RngDistribution

ScatterDimensionNumbersAttr

Attribute that models the dimension information for scatter

Parameters:

Parameter C++ type説明
updateWindowDims ::llvm::ArrayRef<int64_t>寸法
insertedWindowDims ::llvm::ArrayRef<int64_t>寸法
inputBatchingDims ::llvm::ArrayRef<int64_t>寸法
scatterIndicesBatchingDims ::llvm::ArrayRef<int64_t>寸法
scatterDimsToOperandDims ::llvm::ArrayRef<int64_t>寸法
indexVectorDim int64_t

SparsityDescriptorAttr

Describes structured (N:M) sparsity configuration

構文:

#mhlo.sparsity<
  int64_t,   # dimension
  int64_t,   # n
  int64_t   # m
>

This attribute is defined for a sparse dot operation with a structured sparse input tensor. With (N=2,M=4), every 4 consecutive logical elements have exactly 2 non-zero physical elements in the input tensor.

$dimension defines the index of the contracting dimension that is sparse (it has to be the most minor dimension). The additional metadata operand in the sparse dot operation defines which logical elements are zeroed out.

Parameters:

Parameter C++ type説明
寸法int64_t
n int64_t
メートルint64_t

TransposeAttr

Transpose options

構文:

#mhlo.transpose<
  ::mlir::mhlo::Transpose   # value
>

Enum cases:

  • TRANSPOSE_INVALID ( TRANSPOSE_INVALID )
  • NO_TRANSPOSE ( NO_TRANSPOSE )
  • TRANSPOSE ( TRANSPOSE )
  • ADJOINT ( ADJOINT )

Parameters:

Parameter C++ type説明
価値::mlir::mhlo::Transpose an enum of type Transpose

TypeExtensionsAttr

Attribute that extends tensor type with MHLO type properties.

構文:

#mhlo.type_extensions<
  ::llvm::ArrayRef<int64_t>   # bounds
>

This attribute is used to extend MLIR tensor type with MHLO tensor specific properties. These properties aren't modeled in the MLIR type. This attribute is set in the encoding field of the tensor type.

See HLO_BoundedAttrInterface for documentation for bounds .

Parameters:

Parameter C++ type説明
bounds ::llvm::ArrayRef<int64_t>

種類

AsyncBundleType

Opaque collection of other types

構文:

!mhlo.async_bundle<
  ::llvm::ArrayRef<Type>   # types
>

Parameters:

Parameter C++ type説明
種類::llvm::ArrayRef<Type>

Enums

ComparisonDirection

Which comparison operation to perform.

Cases:

シンボル価値
EQ 0 EQ
北東1北東
GE 2 GE
GT 3 GT
4
LT 5 LT

ComparisonType

Which comparison type to use.

Cases:

シンボル価値
NOTYPE 0 NOTYPE
フロート1フロート
TOTALORDER 2 TOTALORDER
SIGNED 3 SIGNED
UNSIGNED 4 UNSIGNED

CustomCallApiVersion

Custom call API version

Cases:

シンボル価値
API_VERSION_UNSPECIFIED 0 API_VERSION_UNSPECIFIED
API_VERSION_ORIGINAL 1 API_VERSION_ORIGINAL
API_VERSION_STATUS_RETURNING 2 API_VERSION_STATUS_RETURNING
API_VERSION_STATUS_RETURNING_UNIFIED 3 API_VERSION_STATUS_RETURNING_UNIFIED
API_VERSION_TYPED_FFI 4 API_VERSION_TYPED_FFI

CustomCallSchedule

Specifies the desired schedule for the custom-call.

Cases:

シンボル価値
なし0なし
最新1最新
EARLIEST 2 EARLIEST

DequantizeMode

Dequantization mode. Only MIN_COMBINED is supported.

Cases:

シンボル価値
MIN_COMBINED 0 MIN_COMBINED

DomainKind

Kind of domain metatdata attached to an HLO domain.

Cases:

シンボル価値
シャーディング0シャーディング

FftType

XLA fast fourier transform type.

Cases:

シンボル価値
FFT 0 FFT
IFFT 1 IFFT
RFFT 2 RFFT
IRFFT 3 IRFFT

FusionKind

fusion kind

Cases:

シンボル価値
kLoop 0 kLoop
kInput 1 kInput
kOutput 2 kOutput
kCustom 3 kCustom

精度

XLA precision for an operand. Has backend specific meaning.

Cases:

シンボル価値
デフォルト0デフォルト
高い1高い
最高2最高
PACKED_NIBBLE 3 PACKED_NIBBLE

RngAlgorithm

XLA PRNG algorithm to be used.

Cases:

シンボル価値
デフォルト0デフォルト
THREE_FRY 1 THREE_FRY
PHILOX 2 PHILOX

RngDistribution

XLA PRNG distribution to be used.

Cases:

シンボル価値
UNIFORM 1 UNIFORM
普通2普通

転置

Transpose options

Cases:

シンボル価値
TRANSPOSE_INVALID 0 TRANSPOSE_INVALID
NO_TRANSPOSE 1 NO_TRANSPOSE
転置2転置
ADJOINT 3 ADJOINT