Opérations
mhlo.abs
(mhlo :: AbsOp)
Opération abdominaux
Syntaxe:
operation ::= `mhlo.abs` $operand attr-dict
`:` custom<SameOperandsAndResultType>(type($operand), type($result))
Effectue une opération abs par élément sur le tenseur operand
et produit un tenseur result
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#abs
Exemple:
%result = mhlo.abs %operand : tensor<3xi32>
Traits : AlwaysSpeculatableImplTrait
, Elementwise
, SameOperandsAndResultShape
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Opérandes :
Opérande | Description |
---|---|
operand | Tenseur classé d'entier sans signe 2/4/8/16/32/64 bits ou de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou Type f8E5M2FNUZ ou type f8E8M0FNU ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou type complexe avec éléments flottants 32 bits ou 64 bits ou uniforme 2/4/8/16/32 bits entier signé quantifié ou quantifié uniforme 2/4/8/16/32 bits par axe entier signé ou Entier non signé quantifié uniforme 2/4/8/16/32 bits ou valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
Résultats:
Résultat | Description |
---|---|
result | Tenseur classé d'entier sans signe 2/4/8/16/32/64 bits ou de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou Type f8E5M2FNUZ ou type f8E8M0FNU ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou entier signé quantifié uniforme 2/4/8/16/32 bits ou 2/4/8/16/32- nombre entier signé uniforme de bits quantifié par axe ou entier non signé quantifié uniforme de 2/4/8/16/32 bits ou valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
mhlo.add
(mhlo :: AddOp)
Ajouter une opération
Syntaxe:
operation ::= `mhlo.add` $lhs `,` $rhs attr-dict
`:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))
Effectue l'addition élément par élément de deux tenseurs lhs
et rhs
et produit un tenseur result
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#add
Exemple:
%result = mhlo.add %lhs, %rhs : tensor<2x2xi32>
Traits : AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Opérandes :
Opérande | Description |
---|---|
lhs | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32 /Entier non signé 64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou 2/4/8/16/32 bits valeurs entières non signées quantifiées uniformes par axe |
rhs | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32 /Entier non signé 64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou 2/4/8/16/32 bits valeurs entières non signées quantifiées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
result | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32 /Entier non signé 64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou 2/4/8/16/32 bits valeurs entières non signées quantifiées uniformes par axe |
mhlo.add_dependency
(mhlo :: AddDependencyOp)
Opération AddDependency
Syntaxe:
operation ::= `mhlo.add_dependency` operands attr-dict `:` functional-type(operands, results)
Cette opération est privée au compilateur XLA, elle n'a donc pas encore de spécification.
De manière informelle, cette opération comporte deux opérandes : un opérande de données et un jeton. Le résultat de l’opération est l’opérande de données. Lorsqu'elle est utilisée avec AfterAll, cette opération permet de commander des opérations sans effets secondaires (celles qui ne produisent pas de valeurs de jeton).
Exemple:
%1 = mhlo.add_dependency %arg0, %0 : (tensor<3x4xf32>, !mhlo.token) -> tensor<3x4xf32>
Traits : AlwaysSpeculatableImplTrait
Interfaces : ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Opérandes :
Opérande | Description |
---|---|
operand | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32 /Entier non signé 64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou valeurs entières non signées quantifiées uniformes de 2/4/8/16/32 bits ou tenseur classé de 2/4/8/16/32 bits quantifié uniforme par axe entier signé ou 2/4/8/16 /32 bits quantifiés uniformément par axe, valeurs entières non signées ou jeton |
token | jeton |
Résultats:
Résultat | Description |
---|---|
output | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32 /Entier non signé 64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou valeurs entières non signées quantifiées uniformes de 2/4/8/16/32 bits ou tenseur classé de 2/4/8/16/32 bits quantifié uniforme par axe entier signé ou 2/4/8/16 /32 bits quantifiés uniformément par axe, valeurs entières non signées ou jeton |
mhlo.after_all
(mhlo::AfterAllOp)
Opération AfterAll
Syntaxe:
operation ::= `mhlo.after_all` $inputs attr-dict
`:` custom<VariadicSameOperandsAndResultType>(ref($inputs), type($inputs), type($result))
Garantit que les opérations produisant les inputs
sont exécutées avant toute opération qui dépend du result
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all
Exemple:
%result = mhlo.after_all %input0, %input1 : !mhlo.token
Traits : AlwaysSpeculatableImplTrait
Interfaces : ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Opérandes :
Opérande | Description |
---|---|
inputs | variadique de jeton |
Résultats:
Résultat | Description |
---|---|
result | jeton |
mhlo.all_gather
(mhlo::AllGatherOp)
Opération AllGather
Au sein de chaque groupe de processus de la grille de processus, concatène les valeurs du tenseur d'opérande de chaque processus le long de all_gather_dim
et produit un tenseur de résultat. Le computation
est appliqué séparément pour chaque opérande dans operands
, produisant un résultat par opérande.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_gather
Exemple:
%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>
Traits : SameOperandsAndResultElementType
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
all_gather_dim | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur n'est pas négative |
replica_groups | ::mlir::DenseIntElementsAttr | Attribut d'éléments entiers sans signe de 64 bits |
channel_handle | ::mlir::mhlo::ChannelHandleAttr | deux entiers de 64 bits « handle » et « type » |
use_global_device_ids | ::mlir::UnitAttr | attribut d'unité |
Opérandes :
Opérande | Description |
---|---|
operands | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
mhlo.all_reduce
(mhlo :: AllReduceOp)
Opération AllReduce
Au sein de chaque groupe de processus dans la grille de processus, applique un computation
de fonction de réduction aux valeurs d'un tenseur d'opérande de chaque processus et produit un tenseur de résultat. Le computation
est appliqué séparément pour chaque opérande dans operands
, produisant un résultat par opérande.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce
Exemple:
%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>
Traits : InferTensorType
, SingleBlockImplicitTerminator<ReturnOp>
, SingleBlock
Interfaces : InferShapedTypeOpInterface
, InferTypeOpInterface
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
replica_groups | ::mlir::DenseIntElementsAttr | Attribut d'éléments entiers sans signe de 64 bits |
channel_handle | ::mlir::mhlo::ChannelHandleAttr | deux entiers de 64 bits « handle » et « type » |
use_global_device_ids | ::mlir::UnitAttr | attribut d'unité |
Opérandes :
Opérande | Description |
---|---|
operands | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
mhlo.all_to_all
(mhlo::AllToAllOp)
Opération AllToAll
Au sein de chaque groupe de processus dans la grille de processus, divise les valeurs du tenseur operand
le long de split_dimension
en parties, disperse les parties divisées entre les processus, concatène les parties dispersées le long concat_dimension
et produit un tenseur result
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_to_all
Exemple:
%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>
Traits : AlwaysSpeculatableImplTrait
, InferTensorType
, SameOperandsElementType
, SameOperandsShape
, SameVariadicOperandSize
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
split_dimension | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur n'est pas négative |
concat_dimension | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur n'est pas négative |
split_count | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur est positive |
replica_groups | ::mlir::DenseIntElementsAttr | Attribut d'éléments entiers sans signe de 64 bits |
channel_handle | ::mlir::mhlo::ChannelHandleAttr | deux entiers de 64 bits « handle » et « type » |
Opérandes :
Opérande | Description |
---|---|
operand | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec des éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe |
mhlo.and
(mhlo::AndOp)
Et le fonctionnement
Syntaxe:
operation ::= `mhlo.and` $lhs `,` $rhs attr-dict
`:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))
Effectue un ET par élément de deux tenseurs lhs
et rhs
et produit un tenseur result
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#and
Exemple:
%result = mhlo.and %lhs, %rhs : tensor<2x2xi32>
Traits : AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Opérandes :
Opérande | Description |
---|---|
lhs | Tenseur classé de pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou valeurs entières non signées 2/4/8/16/32/64 bits |
rhs | Tenseur classé de pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou valeurs entières non signées 2/4/8/16/32/64 bits |
Résultats:
Résultat | Description |
---|---|
result | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32 /Entier non signé 64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou 2/4/8/16/32 bits valeurs entières non signées quantifiées uniformes par axe |
mhlo.async_done
(mhlo :: AsyncDoneOp)
Opération AsyncDone
Cette opération est privée au compilateur XLA, elle n'a donc pas encore de spécification.
De manière informelle, cette opération bloque jusqu'à la fin d'un calcul asynchrone. Il renvoie le résultat final du calcul asynchrone.
Consultez la documentation d'AsyncStart pour plus d'informations.
Interfaces : InferTypeOpInterface
Opérandes :
Opérande | Description |
---|---|
bundle | async_bundle avec n'importe quelle combinaison de tenseurs classés de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou Type f8E8M0FNU ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/ Entier non signé 4/8/16/32/64 bits ou type complexe avec flottant 32 bits ou flottant 64 bits éléments ou entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou quantifié uniforme 2/4/8/16/32 bits par axe signé entier ou quantifié uniforme 2/4/8/16/32 bits par axe, valeurs entières non signées ou valeurs de jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe ou jeton ou tuple imbriqué avec n'importe quelle combinaison de tenseurs classés de Type f4E2M1FN ou type f6E2M3FN ou type f6E3M2FN ou type f8E3M4 ou type f8E4M3 ou type f8E4M3FN ou type f8E4M3FNUZ ou type f8E4M3B11FNUZ ou type f8E5M2 ou type f8E5M2FNUZ ou type f8E8M0FNU ou flottant 16 bits ou 32 bits float ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32/64 bits entier non signé ou type complexe avec des éléments flottants de 32 bits ou 64 bits ou des éléments signés quantifiés uniformes de 2/4/8/16/32 bits entier ou valeurs entières non signées quantifiées uniformes de 2/4/8/16/32 bits ou tenseur classé de 2/4/8/16/32 bits quantifiés uniformes par axe entier signé ou 2/4/8/16/32- bits uniformes quantifiés par axe, valeurs entières non signées ou valeurs de jeton |
mhlo.async_start
(mhlo :: AsyncStartOp)
Opération AsyncStart
Cette opération est privée au compilateur XLA, elle n'a donc pas encore de spécification.
De manière informelle, cette opération lance un calcul asynchrone.
Ceci est utilisé lorsqu'il existe des fonctions qui contiennent à la fois des attentes asynchrones (telles que les DMA) et des calculs sur thread. Par exemple, une fonction peut consister en un calcul, un DMA, un autre calcul, un deuxième DMA et un calcul final. Cela serait représenté par un async_start suivi d'un async_update et d'un async_done. Le async_start effectuerait le premier calcul sur le thread, puis démarrerait le DMA. L'async_update attendrait la fin du DMA si ce n'est pas encore fait, puis exécuterait le deuxième calcul dans la fonction et démarrerait le deuxième DMA. Enfin, async_done attendrait ce dernier DMA, puis exécuterait le dernier calcul qui doit être exécuté sur le thread et renverrait le résultat de ce calcul final.
operands
sont passés au calcul directement called_computation
est la fonction qui sera exécutée de manière asynchrone execution_thread
est le nom du thread dans lequel elle sera exécutée. Le thread principal est appelé « principal ». Tous les sujets ont des noms.
Cela renvoie tout l'état nécessaire entre les opérations asynchrones. Après l'affectation du tampon, les valeurs de retour représentent l'espace nécessaire pour contenir l'entrée, les résultats et tous les blocs-notes nécessaires ou modifiés par l'opération asynchrone.
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
called_computation | ::mlir::FlatSymbolRefAttr | attribut de référence de symbole plat |
execution_thread | ::mlir::StringAttr | attribut de chaîne |
Opérandes :
Opérande | Description |
---|---|
inputs | variadique de tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou 16 bits float ou 32 bits float ou 64 bits float ou bfloat16 type ou pred (AKA booléen ou entier 1 bit) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 /16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits par axe ou Valeurs entières non signées quantifiées uniformes 2/4/8/16/32 bits par axe ou jeton ou tuple imbriqué avec n'importe quelle combinaison de tenseurs classés de Type f4E2M1FN ou type f6E2M3FN ou type f6E3M2FN ou type f8E3M4 ou type f8E4M3 ou type f8E4M3FN ou type f8E4M3FNUZ ou type f8E4M3B11FNUZ ou type f8E5M2 ou type f8E5M2FNUZ ou type f8E8M0FNU ou flottant 16 bits ou 32 bits float ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/4/8/16/32/64 bits entier non signé ou type complexe avec des éléments flottants de 32 bits ou 64 bits ou des éléments signés quantifiés uniformes de 2/4/8/16/32 bits entier ou valeurs entières non signées quantifiées uniformes de 2/4/8/16/32 bits ou tenseur classé de 2/4/8/16/32 bits quantifiés uniformes par axe entier signé ou 2/4/8/16/32- bits uniformes quantifiés par axe, valeurs entières non signées ou valeurs de jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | async_bundle avec n'importe quelle combinaison de tenseurs classés de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou Type f8E8M0FNU ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/ Entier non signé 4/8/16/32/64 bits ou type complexe avec flottant 32 bits ou flottant 64 bits éléments ou entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou quantifié uniforme 2/4/8/16/32 bits par axe signé entier ou quantifié uniforme 2/4/8/16/32 bits par axe, valeurs entières non signées ou valeurs de jeton |
mhlo.async_update
(mhlo :: AsyncUpdateOp)
Opération AsyncUpdate
Cette opération est privée au compilateur XLA, elle n'a donc pas encore de spécification.
De manière informelle, cette opération bloque un calcul asynchrone jusqu'à une barrière de synchronisation. Cela renvoie bundle
après avoir fonctionné dessus.
Consultez la documentation d'AsyncStart pour plus d'informations.
Interfaces : InferTypeOpInterface
Opérandes :
Opérande | Description |
---|---|
bundle | async_bundle avec n'importe quelle combinaison de tenseurs classés de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou Type f8E8M0FNU ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/ Entier non signé 4/8/16/32/64 bits ou type complexe avec flottant 32 bits ou flottant 64 bits éléments ou entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou quantifié uniforme 2/4/8/16/32 bits par axe signé nombre entier ou uniforme 2/4/8/16/32 bits quantifié par axe, valeurs entières non signées ou valeurs de jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | async_bundle avec n'importe quelle combinaison de tenseurs classés de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou Type f8E8M0FNU ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou entier sans signe 2/4/8/16/32/64 bits ou 2/ Entier non signé 4/8/16/32/64 bits ou type complexe avec flottant 32 bits ou flottant 64 bits éléments ou entier signé quantifié uniforme 2/4/8/16/32 bits ou entier non signé quantifié uniforme 2/4/8/16/32 bits ou quantifié uniforme 2/4/8/16/32 bits par axe signé nombre entier ou uniforme 2/4/8/16/32 bits quantifié par axe, valeurs entières non signées ou valeurs de jeton |
mhlo.atan2
(mhlo::Atan2Op)
Opération Atan2
Syntaxe:
operation ::= `mhlo.atan2` $lhs `,` $rhs attr-dict
`:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))
Effectue une opération atan2 par élément sur les tenseurs lhs
et rhs
et produit un tenseur result
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#atan2
Exemple:
%result = mhlo.atan2 %lhs, %rhs : tensor<3xf32>
Traits : AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Opérandes :
Opérande | Description |
---|---|
lhs | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou type float 32 bits ou float 64 bits ou bfloat16 ou type complexe avec éléments float 32 bits ou float 64 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits ou 2/4/8/ Valeurs entières non signées quantifiées uniformes 16/32 bits |
rhs | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou type float 32 bits ou float 64 bits ou bfloat16 ou type complexe avec éléments float 32 bits ou float 64 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits ou 2/4/8/ Valeurs entières non signées quantifiées uniformes 16/32 bits |
Résultats:
Résultat | Description |
---|---|
result | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou type float 32 bits ou float 64 bits ou bfloat16 ou type complexe avec éléments float 32 bits ou float 64 bits ou entier signé quantifié uniforme 2/4/8/16/32 bits ou 2/4/8/ Valeurs entières non signées quantifiées uniformes 16/32 bits |
mhlo.batch_norm_grad
(mhlo :: BatchNormGradOp)
Opération BatchNormGrad
Calcule les gradients de plusieurs entrées de BatchNormTrainingOp rétropropagant à partir de grad_output
et produit les tenseurs grad_operand
, grad_scale
et grad_offset
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_grad
Exemple:
%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>)
Traits : AlwaysSpeculatableImplTrait
, InferTensorType
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
epsilon | ::mlir::FloatAttr | Attribut float 32 bits |
feature_index | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur n'est pas négative |
Opérandes :
Opérande | Description |
---|---|
operand | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou valeurs de type float 32 bits ou float 64 bits ou bfloat16 |
scale | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
mean | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
variance | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
grad_output | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou valeurs de type float 32 bits ou float 64 bits ou bfloat16 |
Résultats:
Résultat | Description |
---|---|
grad_operand | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou valeurs de type float 32 bits ou float 64 bits ou bfloat16 |
grad_scale | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
grad_offset | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
mhlo.batch_norm_inference
(mhlo :: BatchNormInferenceOp)
Opération BatchNormInference
Normalise le tenseur operand
sur toutes les dimensions à l'exception de la dimension feature_index
et produit un tenseur result
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_inference
Exemple:
%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>
Traits : AlwaysSpeculatableImplTrait
, InferTensorType
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
epsilon | ::mlir::FloatAttr | Attribut float 32 bits |
feature_index | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur n'est pas négative |
Opérandes :
Opérande | Description |
---|---|
operand | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou valeurs de type float 32 bits ou float 64 bits ou bfloat16 |
scale | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
offset | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
mean | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
variance | Tenseur 1D de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou Valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16 |
Résultats:
Résultat | Description |
---|---|
result | tenseur classé de type f4E2M1FN ou de type f6E2M3FN ou de type f6E3M2FN ou de type f8E3M4 ou de type f8E4M3 ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E4M3B11FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type f8E8M0FNU ou flottant 16 bits ou valeurs de type float 32 bits ou float 64 bits ou bfloat16 |
mhlo.batch_norm_training
(mhlo :: BatchNormTrainingOp)
Opération BatchNormTraining
Calcule la moyenne et la variance entre les dimensions de lot et spatiales et normalise le tenseur operand
, pour chaque caractéristique de la dimension feature_index
et produit les tenseurs output
, batch_mean
et batch_var
.
Voir : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_training
Exemple:
%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>)
Traits : AlwaysSpeculatableImplTrait
, InferTensorType
Interfaces : ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets : MemoryEffects::Effect{}
Attributs :
Attribut | Type MLIR | Description |
---|---|---|
epsilon | ::mlir::FloatAttr | Attribut float 32 bits |
feature_index | ::mlir::IntegerAttr | Attribut entier sans signe de 64 bits dont la valeur n'est pas négative |
Opérandes :
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou des valeurs de type float 32 bits ou 64 bits ou BFLOAT16 |
scale | Tenseur 1D de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou Type F8E3M4 ou Type F8E4M3 ou F8E4M3FN Type ou type F8E4M2 Float 16 bits ou flotteur 32 bits ou valeurs de type Float ou 64 bits ou BFLOAT16 |
offset | Tenseur 1D de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou Type F8E3M4 ou Type F8E4M3 ou F8E4M3FN Type ou type F8E4M2 Float 16 bits ou flotteur 32 bits ou valeurs de type Float ou 64 bits ou BFLOAT16 |
Résultats:
Résultat | Description |
---|---|
output | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou des valeurs de type float 32 bits ou 64 bits ou BFLOAT16 |
batch_mean | Tenseur 1D de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou Type F8E3M4 ou Type F8E4M3 ou F8E4M3FN Type ou type F8E4M2 Float 16 bits ou flotteur 32 bits ou valeurs de type Float ou 64 bits ou BFLOAT16 |
batch_var | Tenseur 1D de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou Type F8E3M4 ou Type F8E4M3 ou F8E4M3FN Type ou type F8E4M2 Float 16 bits ou flotteur 32 bits ou valeurs de type Float ou 64 bits ou BFLOAT16 |
mhlo.bitcast
(mhlo :: bitcastopop)
Opération de bitcast
Syntaxe:
operation ::= `mhlo.bitcast` operands attr-dict `:` functional-type(operands, results)
Cette opération est privée dans le compilateur XLA, il n'a donc pas encore de spécification.
De manière informelle, cette opération modifie la forme de l'entrée de la manière dont la disposition physique des éléments est inchangée.
Cette opération a besoin d'informations sur disposition pour donner un sens à "l'arrangement physique des éléments", et le support de mise en page dans MHLO est actuellement un travail en cours.
Exemple:
%0 = mhlo.bitcast %arg0 : (tensor<3x4xf32>) -> tensor<3x4x1xf32>
TRAITS: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
mhlo.bitcast_convert
(mhlo :: bitcastConvertop)
Opération BitcastConvert
Syntaxe:
operation ::= `mhlo.bitcast_convert` operands attr-dict `:` functional-type(operands, results)
Effectue une opération Bitcast sur le tenseur operand
et produit un tenseur result
où les bits de l'ensemble du tenseur operand
sont réinterprétés en utilisant le type du tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#bitcast_convert
Exemple:
%result = mhlo.bitcast_convert %operand : (tensor<2xf32>) -> tensor<2x4xi8>
TRAITS: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
mhlo.broadcast
(Mhlo :: Broadcastop)
Opération de diffusion
Cette opération est en train de sortir de Stablehlo, donc elle n'est pas incluse dans la spécification: https://github.com/openxla/stablehlo/issues/3
De manière informelle, cette opération fait la même chose que la diffusion de XLA: https://www.tensorflow.org/xla/operation_semantics#broadcast
Exemple:
%result = mhlo.broadcast %operand, sizes = [1, 2] : (tensor<3xi32>) -> tensor<1x2x3xi32>
TRAITS: AlwaysSpeculatableImplTrait
, InferTensorType
, SameOperandsAndResultElementType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
broadcast_sizes | :: mlir :: dense intelementsAtTr | Attribut d'éléments entiers sans signe 64 bits |
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières uniformes quantifiées par axe |
mhlo.broadcast_in_dim
(Mhlo :: BroadcastIndImop)
Opération de BroadcastIndim
Étend les dimensions et / ou le rang d'un tenseur d'entrée en dupliquant les données dans le tenseur operand
et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#broadcast_in_dim
Exemple:
%result = mhlo.broadcast_in_dim %operand, dims = [2, 1] : (tensor<1x3xi32>) -> tensor<2x3x2xi32>
TRAITS: AlwaysSpeculatableImplTrait
, HLO_CompatibleOperandsAndResultElementType
Interfaces: ConditionallySpeculatable
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
broadcast_dimensions | :: mlir :: dense intelementsAtTr | Attribut d'éléments entiers sans signe 64 bits |
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur de forme statique de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3FN ou Type F8E4M3Fnuz ou F8E4M3B11 Float 16 bits ou flotteur 32 bits ou flotteur 64 bits ou type bfloat16 ou pré (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 / 16/32/64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8/16/32 bits uniformes entier non signé quantisé ou 2/4/8/16/32 bits uniformes quantisés par axe signé entier ou 2/4/8/16/32 bits 2/4/8/16/32 bits uniformes quantifiées par axe Valeurs entières non signées |
mhlo.case
(Mhlo :: Caseop)
Opération de cas
Produit la sortie de l'exécution exactement une function
à partir branches
en fonction de la valeur de index
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#case
Exemple:
%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>)
Traits: RecursiveMemoryEffects
, SingleBlockImplicitTerminator<ReturnOp>
, SingleBlock
Interfaces: InferTypeOpInterface
Opérandes:
Opérande | Description |
---|---|
index | tenseur de valeurs entières sans signe 32 bits |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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 Float 16 bits ou flotteur 32 bits ou flotteur 64 bits ou type bfloat16 ou pré (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 / 16/32/64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8/16/32 bits Valeurs entières non signées quantifiées ou traduits classées de 2/4/8/16/32 bits axe signé entier ou 2/4/8/16/32 bits uniforme quantisé par axe Valeurs entières non signées ou jeton |
mhlo.cbrt
(mhlo :: cbrtop)
Opération CBRT
Syntaxe:
operation ::= `mhlo.cbrt` $operand attr-dict
`:` custom<SameOperandsAndResultType>(type($operand), type($result))
Effectue le fonctionnement de la racine cubique sur l'élément sur le tenseur operand
et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#cbrt
Exemple:
%result = mhlo.cbrt %operand : tensor<4xf32>
TRAITS: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
opérandsandResultType, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou type de complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8 / Valeurs entières uniformes uniformes 16/32 bits |
Résultats:
Résultat | Description |
---|---|
result | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou type de complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8 / Valeurs entières uniformes uniformes 16/32 bits |
mhlo.ceil
(Mhlo :: Ceilop)
Opération de plate-forme
Syntaxe:
operation ::= `mhlo.ceil` $operand attr-dict
`:` custom<SameOperandsAndResultType>(type($operand), type($result))
Effectue le plafond du tenseur d' operand
sur le plan des éléments et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#ceil
Exemple:
%result = mhlo.ceil %operand : tensor<5xf32>
TRAITS: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
opérandsandResultType, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou un flotteur 32 bits ou un flotteur 64 bits ou un type BFLOAT16 ou 2/4/8/16/32 bits uniformes entier signé ou 2/4/8/16/32 bits |
Résultats:
Résultat | Description |
---|---|
result | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou un flotteur 32 bits ou un flotteur 64 bits ou un type BFLOAT16 ou 2/4/8/16/32 bits uniformes entier signé ou 2/4/8/16/32 bits |
mhlo.cholesky
(mhlo :: choleskyop)
Fonctionnement de Cholesky
Calcule la décomposition de Cholesky d'un lot de matrices.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cholesky
Exemple:
%result = mhlo.cholesky %a, lower = true : tensor<3x3xf32>
TRAITS: AlwaysSpeculatableImplTrait
, InferTensorType
, SameOperandsAndResultElementType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
lower | :: mlir :: boolattr | attribut bool |
Opérandes:
Opérande | Description |
---|---|
a | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou Float 32 bits ou Float 64 bits ou BFLOAT16 Type ou Type de complexe avec des valeurs d'éléments flottants de flotteur 32 bits ou 64 bits |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou Float 32 bits ou Float 64 bits ou BFLOAT16 Type ou Type de complexe avec des valeurs d'éléments flottants de flotteur 32 bits ou 64 bits |
mhlo.clamp
(Mhlo :: Clampop)
Opération de pince
Syntaxe:
operation ::= `mhlo.clamp` $min `,` $operand `,` $max attr-dict
`:` custom<SameOperandsAndResultType>(type($min), type($operand), type($max), type($result))
Graque chaque élément du tenseur operand
entre une valeur minimale et maximale et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#clamp
Exemple:
%result = mhlo.clamp %min, %operand, %max : tensor<3xi32>
TRAITS: AlwaysSpeculatableImplTrait
SPECULATableImpltrait, HLO_BroadcastingElementwise
, InferTensorType
, SameOperandsAndResultElementType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Opérandes:
Opérande | Description |
---|---|
min | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
max | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
result | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
mhlo.collective_broadcast
(Mhlo :: CollectiveBroadcastop)
Opération CollectiveBroadcast
Dans chaque groupe de processus de la grille de processus, envoyez la valeur du tenseur operand
du processus source aux processus cibles et produisez un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#Collective_broadcast
Exemple:
%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>
Traits: CompatibleOperandsAndResultType
Interfaces: InferShapedTypeOpInterface
, InferTypeOpInterface
Attributs:
Attribut | Type mlir | Description |
---|---|---|
replica_groups | :: mlir :: dense intelementsAtTr | Attribut d'éléments entiers sans signe 64 bits |
channel_handle | :: Mlir :: Mhlo :: ChannelHandleattr | Deux entiers 64 bits «manche» et «type» |
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
mhlo.collective_permute
(Mhlo :: CollectivePermuteop)
Opération collectivePermute
Dans chaque groupe de processus de la grille de processus, envoie la valeur du tenseur operand
du processus source au processus cible et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_permute
Exemple:
%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>
TRAITS: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
source_target_pairs | :: mlir :: dense intelementsAtTr | Attribut d'éléments entiers sans signe 64 bits |
channel_handle | :: Mlir :: Mhlo :: ChannelHandleattr | Deux entiers 64 bits «manche» et «type» |
Opérandes:
Opérande | Description |
---|---|
operand | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
mhlo.compare
(mhlo :: compareop)
Comparez le fonctionnement
Syntaxe:
operation ::= `mhlo.compare` $comparison_direction `,` $lhs `,` $rhs (`,` $compare_type^)?
attr-dict `:` functional-type(operands, results)
Effectue une comparaison par élément des tenseurs lhs
et rhs
en fonction de comparison_direction
et compare_type
, et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#compare
Exemple:
%result = mhlo.compare LT, %lhs, %rhs, FLOAT : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
TRAITS: AlwaysSpeculatableImplTrait
SPECULATableImpltrait, Elementwise
, InferTensorType
, SameOperandsAndResultShape
, SameOperandsElementType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
comparison_direction | :: Mlir :: Mhlo :: ComparisonDirectionAtr | Quelle opération de comparaison à effectuer. |
compare_type | :: mlir :: mhlo :: comparaisonypeattr | Quel type de comparaison à utiliser. |
Opérandes:
Opérande | Description |
---|---|
lhs | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
rhs | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé des valeurs de préd (aka boolean ou 1 bits) |
mhlo.complex
(MHLO :: Complexop)
Opération complexe
Syntaxe:
operation ::= `mhlo.complex` operands attr-dict
`:` custom<ComplexOpType>(type($lhs), type($rhs), type($result))
Effectue une conversion d'élément en une valeur complexe à partir d'une paire de valeurs réelles et imaginaires, lhs
et rhs
, et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#complex
Exemple:
%result = mhlo.complex %lhs, %rhs : tensor<2xcomplex<f32>>
TRAITS: AlwaysSpeculatableImplTrait
, Elementwise
, SameOperandsAndResultShape
, SameOperandsElementType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Opérandes:
Opérande | Description |
---|---|
lhs | Tensor classé de Float 32 bits ou de valeurs de flotteur 64 bits |
rhs | Tensor classé de Float 32 bits ou de valeurs de flotteur 64 bits |
Résultats:
Résultat | Description |
---|---|
result | Tensor classé de type complexe avec des valeurs de flotteur 32 bits ou des éléments flottants 64 bits |
mhlo.composite
(mhlo :: compositeop)
Opération composite
Syntaxe:
operation ::= `mhlo.composite` $name $inputs attr-dict `:` functional-type(operands, results)
Résume une opération constituée (composée) d'autres opérations stablehlo, en prenant inputs
et composite_attributes
et en produisant results
. La sémantique de l'OP est mise en œuvre par l'attribut decomposition
. L'OP composite
peut être remplacé par sa décomposition sans changer la sémantique de programme. Dans les cas où l'inclinaison de la décomposition ne fournit pas la même sémantique OP, préférez utiliser custom_call
.
Le champ version
(par défaut à 0
) est utilisé pour indiquer le changement de sémantique d'un composite.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#composite
Exemple:
%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>
Interfaces: SymbolUserOpInterface
Attributs:
Attribut | Type mlir | Description |
---|---|---|
name | :: mlir :: stringattr | attribut de chaîne |
composite_attributes | :: mlir :: dictionaryattr | Dictionnaire des valeurs d'attribut nommées |
decomposition | :: Mlir :: FlatsymbolRefattr | Attribut de référence à symbole plat |
version | :: Mlir :: Integerattr | Attribut entier sans signe 32 bits |
Opérandes:
Opérande | Description |
---|---|
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 Float 16 bits ou flotteur 32 bits ou flotteur 64 bits ou type bfloat16 ou pré (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 / 16/32/64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8/16/32 bits uniformes entier non signé quantisé ou 2/4/8/16/32 bits uniformes quantisés par axe signé entier ou 2/4/8/16/32 bits 2/4/8/16/32 bits uniformes quantifiées par axe Valeurs entières non signées ou tuple de jeton ou imbriqué avec toute combinaison de tenseur classé de F4E2M1FN Type ou F6E2M3FN Type ou type F6E3M2FN ou type F8E3M4 ou type F8E4M3 ou F8E4M3B11FNUZ Type ou F8E5M2 float ou 64 bits float ou bfloat16 type ou pred (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 2/4/8/16/32 bits uniformes quantifiés signés entier ou 2/4/8/16/32 bits uniformes de valeurs entières non signées quantifiées ou tenseur classé de 2/4/8/16/32 bits quantifié par axe entier signé ou 2/4/8/16/32- bit uniforme quantifié par axe Valeurs entières non signées ou valeurs de jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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 Float 16 bits ou flotteur 32 bits ou flotteur 64 bits ou type bfloat16 ou pré (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 / 16/32/64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8/16/32 bits uniformes entier non signé quantisé ou 2/4/8/16/32 bits uniformes quantisés par axe signé entier ou 2/4/8/16/32 bits 2/4/8/16/32 bits uniformes quantifiées par axe Valeurs entières non signées ou tuple de jeton ou imbriqué avec toute combinaison de tenseur classé de F4E2M1FN Type ou F6E2M3FN Type ou type F6E3M2FN ou type F8E3M4 ou type F8E4M3 ou F8E4M3B11FNUZ Type ou F8E5M2 float ou 64 bits float ou bfloat16 type ou pred (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 2/4/8/16/32 bits uniformes quantifiés signés entier ou 2/4/8/16/32 bits uniformes de valeurs entières non signées quantifiées ou tenseur classé de 2/4/8/16/32 bits quantifié par axe entier signé ou 2/4/8/16/32- bit uniforme quantifié par axe Valeurs entières non signées ou valeurs de jeton |
mhlo.concatenate
(mhlo :: concatenateop)
Opération de concaténate
COMPATENNES UN NOMBRE VARIADIQUE DE TENSEURS DANS inputs
le long de la dimension dimension
dans le même ordre que les arguments donnés et produit un tenseur result
.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#concatenate
Exemple:
%result = mhlo.concatenate %input0, %input1, dim = 0 : (tensor<3x2xi64>, tensor<1x2xi64>) -> tensor<4x2xi64>
TRAITS: AlwaysSpeculatableImplTrait
, SameOperandsAndResultElementType
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
dimension | :: Mlir :: Integerattr | Attribut entier 64 bits sans signe dont la valeur n'est pas négative |
Opérandes:
Opérande | Description |
---|---|
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 Float 16 bits ou flotteur 32 bits ou flotteur 64 bits ou type bfloat16 ou pré (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8 / 16/32/64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé uniforme ou 2/4/8/16/32 bits uniformes entier non signé quantisé ou 2/4/8/16/32 bits uniformes quantisés par axe signé entier ou 2/4/8/16/32 bits 2/4/8/16/32 bits uniformes quantifiées par axe Valeurs entières non signées |
Résultats:
Résultat | Description |
---|---|
"anonyme" | Tenseur classé de type F4E2M1FN ou type F6E2M3FN ou type F6E3M2FN ou type F8E3M4 ou Type F8E4M3 ou F8E4M3B11fnUz Type ou Type F8e5M2 ou flotteur 32 bits ou flotteur 64 bits ou type BFLOAT16 ou préd (aka boolean ou entier 1 bits) ou 2/4/8/16/32/64 bits entier sans signe ou 2/4/8/16/32 / 64 bits entier non signé ou type complexe avec un flotteur 32 bits ou des éléments flottants 64 bits ou 2/4/8/16/32 bits entier signé quantifié uniforme ou 2/4/8/16/32 bits entier uniforme uniforme unise Valeurs entières non signées uniformes par axe |
mhlo.constant
(Mhlo :: Constantop)
Opération constante
Produit un tenseur output
à partir d'une value
constante.
Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#constant
Exemple:
%output = mhlo.constant dense<[[0.0, 1.0], [2.0, 3.0]]> : tensor<2x2xf32>
Traits: AlwaysSpeculatableImplTrait
, ConstantLike
Interfaces: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effets: MemoryEffects::Effect{}
Attributs:
Attribut | Type mlir | Description |
---|---|---|
value | :: mlir :: elementSattr | Attribut vectoriel / tensor constant |
Résultats:
Résultat | Description |
---|---|
output | 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.convert
(mhlo::ConvertOp)
Convert operation
Syntaxe:
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
Exemple:
%result = mhlo.convert %operand : (tensor<3xi32>) -> tensor<3xcomplex<f32>>
Traits: AlwaysSpeculatableImplTrait
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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.
Exemple:
%0 = mhlo.copy %arg0 : tensor<f32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
cross_program_prefetch_index | ::mlir::IntegerAttr | 32-bit signless integer attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.cosine %operand : tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.count_leading_zeros %operand : tensor<2x2xi8>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
operand | ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%output = mhlo.create_token : !mhlo.token
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Résultats:
Résultat | Description |
---|---|
output | jeton |
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
replica_groups | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.divide %lhs, %rhs : tensor<4xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%0 = mhlo.dot %arg0, %arg1 : (tensor<1x2xi32>, tensor<2x1xi32>) -> tensor<1x1xi32>
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
precision_config | ::mlir::ArrayAttr | Precision Config attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dimension_numbers | ::mlir::mhlo::GatherDimensionNumbersAttr | Attribute that models the dimension information for gather |
indices_are_sorted | ::mlir::BoolAttr | bool attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%0 = mhlo.dynamic_iota %arg0, dim = 0 : (tensor<1xindex>) -> tensor<4xi32>
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
iota_dimension | ::mlir::IntegerAttr | 64-bit signless integer attribute whose value is non-negative |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
slice_sizes | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
einsum_config | ::mlir::StringAttr | string attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.erf %operand : tensor<2x2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.exponential %operand : tensor<2x2xf64>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.exponential_minus_one %operand : tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
fft_type | ::mlir::mhlo::FftTypeAttr | XLA fast fourier transform type. |
fft_length | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.floor %operand : tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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:
Attribut | MLIR Type | Description |
---|---|---|
fusion_kind | ::mlir::mhlo::FusionKindAttr | fusion kind |
output_operand_aliases | ::mlir::ArrayAttr | Aliasing attribute for outputs and operands of Fusion |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | 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 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dimension | ::mlir::IntegerAttr | 64-bit signless integer attribute whose value is non-negative |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | tensor of 32-bit signless integer values |
mhlo.get_tuple_element
(mhlo::GetTupleElementOp)
GetTupleElement operation
Syntaxe:
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
index | ::mlir::IntegerAttr | 32-bit signless integer attribute whose value is non-negative |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Traits: RecursiveMemoryEffects
, SingleBlockImplicitTerminator<ReturnOp>
, SingleBlock
Interfaces: InferTypeOpInterface
Operands:
Operand | Description |
---|---|
pred | ranked tensor of pred (AKA boolean or 1-bit integer) values |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.imag %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%results:2 = "mhlo.infeed"(%token) {
infeed_config = ""
} : (!mhlo.token) -> (tensor<3x3x3xi32>, !mhlo.token)
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
infeed_config | ::mlir::StringAttr | string attribute |
layout | ::mlir::ArrayAttr | array attribute |
Operands:
Operand | Description |
---|---|
token | jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%output = mhlo.iota dim = 0 : tensor<4x5xi32>
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
iota_dimension | ::mlir::IntegerAttr | 64-bit signless integer attribute whose value is non-negative |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%y = mhlo.is_finite %x : (tensor<7xf32>) -> tensor<7xi1>
Traits: AlwaysSpeculatableImplTrait
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
y | ranked tensor of pred (AKA boolean or 1-bit integer) values |
mhlo.log
(mhlo::LogOp)
Log operation
Syntaxe:
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
Exemple:
%result = mhlo.log %operand : tensor<2x2xf64>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.log_plus_one %operand : tensor<6xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.logistic %operand : tensor<2x2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dimensions | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.maximum %lhs, %rhs : tensor<4xf32>
Traits: AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.minimum %lhs, %rhs : tensor<4xf32>
Traits: AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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:
Operand | Description |
---|---|
shapes | variadic of 1D tensor of index values |
Résultats:
Résultat | Description |
---|---|
results | variadic of 1D tensor of index values |
mhlo.multiply
(mhlo::MulOp)
Mul operation
Syntaxe:
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
Exemple:
%result = mhlo.multiply %lhs, %rhs : tensor<2xi32>
Traits: AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.negate %operand : tensor<2x3xi32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.not %operand : tensor<5x3x1xi1>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.or %lhs, %rhs : tensor<2xi1>
Traits: AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%result = "mhlo.outfeed"(%input0, %token) {
outfeed_config = ""
} : (tensor<3x3x3xi32>, !mhlo.token) -> !mhlo.token
Interfaces: InferTypeOpInterface
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
outfeed_config | ::mlir::StringAttr | string attribute |
Operands:
Operand | Description |
---|---|
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 | jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | jeton |
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.partition_id : tensor<ui32>
Interfaces: InferTypeOpInterface
Résultats:
Résultat | Description |
---|---|
"anonyme" | ranked tensor of 32-bit unsigned integer values |
mhlo.popcnt
(mhlo::PopulationCountOp)
PopulationCount operation
Syntaxe:
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
Exemple:
%result = mhlo.popcnt %operand : tensor<4xi8>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
operand | ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.power %lhs, %rhs : tensor<6xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.real %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
channel_handle | ::mlir::mhlo::ChannelHandleAttr | two 64-bit integers 'handle' and 'type' |
is_host_transfer | ::mlir::BoolAttr | bool attribute |
Operands:
Operand | Description |
---|---|
token | jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dimensions | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%output = mhlo.reduce_precision %operand, format = e5m2 : tensor<6xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.remainder %lhs, %rhs : tensor<4xi64>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.replica_id : tensor<ui32>
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Résultats:
Résultat | Description |
---|---|
"anonyme" | ranked tensor of 32-bit unsigned integer values |
mhlo.reshape
(mhlo::ReshapeOp)
Reshape operation
Syntaxe:
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
Exemple:
%result = mhlo.reshape %operand : (tensor<2xf32>) -> tensor<1x2xf32>
Traits: AlwaysSpeculatableImplTrait
, HLO_CompatibleOperandsAndResultElementType
Interfaces: ConditionallySpeculatable
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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:
Attribut | MLIR Type | Description |
---|---|---|
dimensions | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%result = mhlo.rng %a, %b, %shape, distribution = NORMAL : (tensor<i32>, tensor<i32>, tensor<2xi64>) -> tensor<3x3xi32>
Traits: InferTensorType
Interfaces: InferShapedTypeOpInterface
, InferTypeOpInterface
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
rng_distribution | ::mlir::mhlo::RngDistributionAttr | XLA PRNG distribution to be used. |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
rng_algorithm | ::mlir::mhlo::RngAlgorithmAttr | XLA PRNG algorithm to be used. |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.round_nearest_afz %operand : tensor<5xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.round_nearest_even %operand : tensor<5xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.rsqrt %operand : tensor<2x2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
channel_handle | ::mlir::mhlo::ChannelHandleAttr | two 64-bit integers 'handle' and 'type' |
is_host_transfer | ::mlir::BoolAttr | bool attribute |
Operands:
Operand | Description |
---|---|
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 | jeton |
Résultats:
Résultat | Description |
---|---|
"anonyme" | jeton |
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dimension | ::mlir::IntegerAttr | 64-bit signless integer attribute whose value is non-negative |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.shift_left %lhs, %rhs : tensor<6xi8>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.shift_right_arithmetic %lhs, %rhs : tensor<6xi8>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.shift_right_logical %lhs, %rhs : tensor<6xi8>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.sign %operand : tensor<7xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.sine %operand : tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dimension | ::mlir::IntegerAttr | 64-bit signless integer attribute |
is_stable | ::mlir::BoolAttr | bool attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.sqrt %operand : tensor<2x2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.subtract %lhs, %rhs : tensor<2xi32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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.
Exemple:
%0 = mhlo.tan %arg0 : tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%result = mhlo.tanh %operand : tensor<2xf32>
Traits: AlwaysSpeculatableImplTrait
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%values, %indices = mhlo.topk(%operand, k=5, largest=true)
: tensor<100xf32> -> (tensor<5xf32>, tensor<5xi32>)
Traits: InferTensorType
, RecursiveMemoryEffects
Interfaces: InferShapedTypeOpInterface
, InferTypeOpInterface
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
k | ::mlir::IntegerAttr | 64-bit signless integer attribute |
largest | ::mlir::BoolAttr | bool attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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.
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
dim | ::mlir::IntegerAttr | 64-bit signless integer attribute |
batch_dims | ::mlir::IntegerAttr | 64-bit signless integer attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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.
Exemple:
mhlo.trace %arg0, "In test code." : tensor<5x1x5xi32>
Attributes:
Attribut | MLIR Type | Description |
---|---|---|
tag | ::mlir::StringAttr | string attribute |
Operands:
Operand | Description |
---|---|
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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
permutation | ::mlir::DenseIntElementsAttr | 64-bit signless integer elements attribute |
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Exemple:
%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:
Attribut | MLIR Type | Description |
---|---|---|
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:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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
Exemple:
%result = mhlo.tuple %val0, %val1 : tuple<tensor<2xf32>, tuple<tensor<i32>>>
Traits: AlwaysSpeculatableImplTrait
Interfaces: ConditionallySpeculatable
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Syntaxe:
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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
Exemple:
%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 | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
"anonyme" | 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
Syntaxe:
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:
Attribut | MLIR Type | Description |
---|---|---|
delta | ::mlir::IntegerAttr | 64-bit signless integer attribute |
Résultats:
Résultat | Description |
---|---|
"anonyme" | statically shaped tensor of 64-bit unsigned integer values |
mhlo.xor
(mhlo::XorOp)
Xor operation
Syntaxe:
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
Exemple:
%result = mhlo.xor %lhs, %rhs : tensor<2xi32>
Traits: AlwaysSpeculatableImplTrait
, Commutative
, CompatibleOperandsAndResultType
, Elementwise
, SameOperandsAndResultShape
Interfaces: ConditionallySpeculatable
, InferShapedTypeOpInterface
, InferTypeOpInterface
, NoMemoryEffect (MemoryEffectOpInterface)
Effects: MemoryEffects::Effect{}
Operands:
Operand | Description |
---|---|
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 |
Résultats:
Résultat | Description |
---|---|
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 |
Attributs
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 ...
}
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
argTupleIndices | ::llvm::ArrayRef<int64_t> | Dimension |
resultIndex | int64_t | |
resultTupleIndices | ::llvm::ArrayRef<int64_t> | Dimension |
isMustAlias | bool |
ChannelHandleAttr
two 64-bit integers 'handle' and 'type'
Syntaxe:
#mhlo.channel_handle<
int64_t, # handle
int64_t # type
>
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
poignée | int64_t | |
taper | int64_t |
ComparisonDirectionAttr
Which comparison operation to perform.
Syntaxe:
#mhlo.comparison_direction<
::mlir::mhlo::ComparisonDirection # value
>
Enum cases:
- EQ (
EQ
) - NE (
NE
) - GE (
GE
) - GT (
GT
) - LE (
LE
) - LT (
LT
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::ComparisonDirection | an enum of type ComparisonDirection |
ComparisonTypeAttr
Which comparison type to use.
Syntaxe:
#mhlo.comparison_type<
::mlir::mhlo::ComparisonType # value
>
Enum cases:
- NOTYPE (
NOTYPE
) - FLOAT (
FLOAT
) - TOTALORDER (
TOTALORDER
) - SIGNED (
SIGNED
) - UNSIGNED (
UNSIGNED
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::ComparisonType | an enum of type ComparisonType |
ConvDimensionNumbersAttr
Structure of dimension information for conv op
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
inputBatchDimension | int64_t | |
inputFeatureDimension | int64_t | |
inputSpatialDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
kernelInputFeatureDimension | int64_t | |
kernelOutputFeatureDimension | int64_t | |
kernelSpatialDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
outputBatchDimension | int64_t | |
outputFeatureDimension | int64_t | |
outputSpatialDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
CrossProgramPrefetchAttr
Argument that is prefetched from another program
Syntaxe:
#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.
Par exemple,
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.
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
paramètre | int64_t | |
indices | ::llvm::ArrayRef<int64_t> | Dimension |
compenser | std::optional<int64_t> |
CustomCallScheduleAttr
Specifies the desired schedule for the custom-call.
Syntaxe:
#mhlo.custom_call_schedule<
::mlir::mhlo::CustomCallSchedule # value
>
Enum cases:
- NONE (
NONE
) - LATEST (
LATEST
) - EARLIEST (
EARLIEST
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::CustomCallSchedule | an enum of type CustomCallSchedule |
DequantizeModeAttr
Dequantization mode. Only MIN_COMBINED is supported.
Syntaxe:
#mhlo.dequantize_mode<
::mlir::mhlo::DequantizeMode # value
>
Enum cases:
- MIN_COMBINED (
MIN_COMBINED
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::DequantizeMode | an enum of type DequantizeMode |
DomainKindAttr
Kind of domain metatdata attached to an HLO domain.
Syntaxe:
#mhlo.kind<
::mlir::mhlo::DomainKind # value
>
Enum cases:
- sharding (
sharding
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::DomainKind | an enum of type DomainKind |
DotAlgorithmAttr
Attribute that models the algorithm constraints to use for computing dot.
Syntaxe:
#mhlo.dot_algorithm<
Type, # lhsPrecisionType
Type, # rhsPrecisionType
Type, # accumulationType
int64_t, # lhsComponentCount
int64_t, # rhsComponentCount
int64_t, # numPrimitiveOperations
bool # allowImpreciseAccumulation
>
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
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.
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
lhsBatchingDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
rhsBatchingDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
lhsContractingDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
rhsContractingDimensions | ::llvm::ArrayRef<int64_t> | Dimension |
FftTypeAttr
XLA fast fourier transform type.
Syntaxe:
#mhlo.fft_type<
::mlir::mhlo::FftType # value
>
Enum cases:
- FFT (
FFT
) - IFFT (
IFFT
) - RFFT (
RFFT
) - IRFFT (
IRFFT
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::FftType | an enum of type FftType |
FusionKindAttr
fusion kind
Syntaxe:
#mhlo.fusion_kind<
::mlir::mhlo::FusionKind # value
>
Enum cases:
- kLoop (
kLoop
) - kInput (
kInput
) - kOutput (
kOutput
) - kCustom (
kCustom
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::FusionKind | an enum of type FusionKind |
GatherDimensionNumbersAttr
Attribute that models the dimension information for gather
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
offsetDims | ::llvm::ArrayRef<int64_t> | Dimension |
collapsedSliceDims | ::llvm::ArrayRef<int64_t> | Dimension |
operandBatchingDims | ::llvm::ArrayRef<int64_t> | Dimension |
startIndicesBatchingDims | ::llvm::ArrayRef<int64_t> | Dimension |
startIndexMap | ::llvm::ArrayRef<int64_t> | Dimension |
indexVectorDim | int64_t |
OutputOperandAliasAttr
Attribute that models the alias relationship of output and operand of a CustomCall op
Syntaxe:
#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>.
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
outputTupleIndices | ::llvm::ArrayRef<int64_t> | Dimension |
operandIndex | int64_t | |
operandTupleIndices | ::llvm::ArrayRef<int64_t> | Dimension |
PrecisionAttr
XLA precision for an operand. Has backend specific meaning.
Syntaxe:
#mhlo.precision<
::mlir::mhlo::Precision # value
>
Enum cases:
- DEFAULT (
DEFAULT
) - HIGH (
HIGH
) - HIGHEST (
HIGHEST
) - PACKED_NIBBLE (
PACKED_NIBBLE
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::Precision | an enum of type Precision |
RngAlgorithmAttr
XLA PRNG algorithm to be used.
Syntaxe:
#mhlo.rng_algorithm<
::mlir::mhlo::RngAlgorithm # value
>
Enum cases:
- DEFAULT (
DEFAULT
) - THREE_FRY (
THREE_FRY
) - PHILOX (
PHILOX
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::RngAlgorithm | an enum of type RngAlgorithm |
RngDistributionAttr
XLA PRNG distribution to be used.
Syntaxe:
#mhlo.rng_distribution<
::mlir::mhlo::RngDistribution # value
>
Enum cases:
- UNIFORM (
UNIFORM
) - NORMAL (
NORMAL
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::RngDistribution | an enum of type RngDistribution |
ScatterDimensionNumbersAttr
Attribute that models the dimension information for scatter
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
updateWindowDims | ::llvm::ArrayRef<int64_t> | Dimension |
insertedWindowDims | ::llvm::ArrayRef<int64_t> | Dimension |
inputBatchingDims | ::llvm::ArrayRef<int64_t> | Dimension |
scatterIndicesBatchingDims | ::llvm::ArrayRef<int64_t> | Dimension |
scatterDimsToOperandDims | ::llvm::ArrayRef<int64_t> | Dimension |
indexVectorDim | int64_t |
SparsityDescriptorAttr
Describes structured (N:M) sparsity configuration
Syntaxe:
#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.
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
dimension | int64_t | |
n | int64_t | |
m | int64_t |
TransposeAttr
Transpose options
Syntaxe:
#mhlo.transpose<
::mlir::mhlo::Transpose # value
>
Enum cases:
- TRANSPOSE_INVALID (
TRANSPOSE_INVALID
) - NO_TRANSPOSE (
NO_TRANSPOSE
) - TRANSPOSE (
TRANSPOSE
) - ADJOINT (
ADJOINT
)
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
valeur | ::mlir::mhlo::Transpose | an enum of type Transpose |
TypeExtensionsAttr
Attribute that extends tensor type with MHLO type properties.
Syntaxe:
#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
.
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
bornes | ::llvm::ArrayRef<int64_t> |
Espèces
AsyncBundleType
Opaque collection of other types
Syntaxe:
!mhlo.async_bundle<
::llvm::ArrayRef<Type> # types
>
Paramètres :
Paramètre | C++ type | Description |
---|---|---|
genres | ::llvm::ArrayRef<Type> |
Enums
ComparisonDirection
Which comparison operation to perform.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
égaliseur | 0 | égaliseur |
NE | 1 | NE |
GE | 2 | GE |
GT | 3 | GT |
LE | 4 | LE |
LT | 5 | LT |
ComparisonType
Which comparison type to use.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
NOTYPE | 0 | NOTYPE |
FLOTTER | 1 | FLOTTER |
TOTALORDER | 2 | TOTALORDER |
SIGNÉ | 3 | SIGNÉ |
UNSIGNED | 4 | UNSIGNED |
CustomCallApiVersion
Custom call API version
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
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.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
AUCUN | 0 | AUCUN |
DERNIER | 1 | DERNIER |
EARLIEST | 2 | EARLIEST |
DequantizeMode
Dequantization mode. Only MIN_COMBINED is supported.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
MIN_COMBINED | 0 | MIN_COMBINED |
DomainKind
Kind of domain metatdata attached to an HLO domain.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
sharding | 0 | sharding |
FftType
XLA fast fourier transform type.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
FFT | 0 | FFT |
IFFT | 1 | IFFT |
RFFT | 2 | RFFT |
IRFFT | 3 | IRFFT |
FusionKind
fusion kind
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
kLoop | 0 | kLoop |
kInput | 1 | kInput |
kOutput | 2 | kOutput |
kCustom | 3 | kCustom |
Précision
XLA precision for an operand. Has backend specific meaning.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
DÉFAUT | 0 | DÉFAUT |
HAUT | 1 | HAUT |
HIGHEST | 2 | HIGHEST |
PACKED_NIBBLE | 3 | PACKED_NIBBLE |
RngAlgorithm
XLA PRNG algorithm to be used.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
DÉFAUT | 0 | DÉFAUT |
THREE_FRY | 1 | THREE_FRY |
PHILOX | 2 | PHILOX |
RngDistribution
XLA PRNG distribution to be used.
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
UNIFORME | 1 | UNIFORME |
NORMALE | 2 | NORMALE |
Transposer
Transpose options
Cas :
Symbole | Valeur | Chaîne |
---|---|---|
TRANSPOSE_INVALID | 0 | TRANSPOSE_INVALID |
NO_TRANSPOSE | 1 | NO_TRANSPOSE |
TRANSPOSER | 2 | TRANSPOSER |
ADJOINT | 3 | ADJOINT |