'mhlo' Dialecte

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 4/8/16/32/64 bits ou de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type float 16 bits ou de type float 32 bits ou de type float 64 bits ou de type bfloat16 ou type complexe avec des éléments flottants de 32 bits ou de 64 bits ou un entier signé quantifié uniforme de 4/8/16/32 bits ou un entier signé quantifié uniforme de 4/8/16/32 bits par axe ou un entier signé de 4/8/16/ Entier non signé quantifié uniforme de 32 bits ou valeurs entières non signées quantifiées uniformes de 4/8/16/32 bits par axe

Résultats:

Résultat Description
result Tenseur classé d'entier sans signe 4/8/16/32/64 bits ou de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de type float 16 bits ou de type float 32 bits ou de type float 64 bits ou de type bfloat16 ou Entier signé quantifié uniforme 4/8/16/32 bits ou Entier signé quantifié uniforme 4/8/16/32 bits par axe ou Entier non signé quantifié uniforme 4/8/16/32 bits ou 4/8/16/ Valeurs entières non signées quantifiées uniformes sur 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits par axe
rhs Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits par axe

Résultats:

Résultat Description
result Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou 4/8/16/32 bits quantifiés uniformes valeurs entières non signées ou tenseur classé de 4/8/16/32 bits quantifiés uniformes par axe entier signé ou 4/8/16/32 bits quantifiés uniformes par axe valeurs entières non signées ou jeton
token jeton

Résultats:

Résultat Description
output Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou 4/8/16/32 bits quantifiés uniformes valeurs entières non signées ou tenseur classé de 4/8/16/32 bits quantifiés uniformes par axe entier signé ou 4/8/16/32 bits quantifiés uniformes 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

Les attributs:

Attribut Type MLIR Description
all_gather_dim ::mlir::IntegerAttr Attribut entier sans signe de 64 bits
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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe signé entier ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs

Résultats:

Résultat Description
"anonyme" variadique de tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe signé entier ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs

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

Les 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs

Résultats:

Résultat Description
"anonyme" variadique de tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs

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

Les attributs:

Attribut Type MLIR Description
split_dimension ::mlir::IntegerAttr Attribut entier sans signe de 64 bits
concat_dimension ::mlir::IntegerAttr Attribut entier sans signe de 64 bits
split_count ::mlir::IntegerAttr Attribut entier sans signe de 64 bits
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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs

Résultats:

Résultat Description
"anonyme" variadique de tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs

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 4/8/16/32/64 bits ou valeurs entières non signées 4/8/16/32/64 bits
rhs Tenseur classé de pred (AKA booléen ou entier 1 bit) ou entier sans signe 4/8/16/32/64 bits ou valeurs entières non signées 4/8/16/32/64 bits

Résultats:

Résultat Description
result Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits 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

Les 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
bundle async_bundle avec n'importe quelle combinaison de tenseurs classés de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 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 4/8/16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/ Entier signé quantifié uniforme de 32 bits ou entier non signé quantifié uniforme de 4/8/16/32 bits ou quantifié uniforme de 4/8/16/32 bits par axe entier signé ou quantifié uniforme de 4/8/16/32 bits par valeurs entières non signées de l'axe ou valeurs de jeton

Résultats:

Résultat Description
"anonyme" variadique de tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs ou jeton ou tuple imbriqué avec toute combinaison de tenseurs classés de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA boolean ou entier 1 bit) ou entier sans signe 4/8/16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou 4 /8/16/32 bits nombre entier signé quantifié uniforme ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits ou tenseur classé de 4/8/16/32 bits quantifié uniforme par axe entier signé ou 4/8 /16/32 bits quantifiés uniformes 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. async_update attendrait la fin du DMA si ce n'était 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.

Les 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4 /8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/32 bits entier signé quantifié uniforme ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou quantifié uniforme 4/8/16/32 bits par axe entier non signé valeurs ou jeton ou tuple imbriqué avec toute combinaison de tenseurs classés de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA boolean ou entier 1 bit) ou entier sans signe 4/8/16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou 4 /8/16/32 bits nombre entier signé quantifié uniforme ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits ou tenseur classé de 4/8/16/32 bits quantifié uniforme par axe entier signé ou 4/8 /16/32 bits quantifiés uniformes 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 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 4/8/16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/ Entier signé quantifié uniforme de 32 bits ou entier non signé quantifié uniforme de 4/8/16/32 bits ou quantifié uniforme de 4/8/16/32 bits par axe entier signé ou quantifié uniforme de 4/8/16/32 bits par valeurs entières non signées de l'axe 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

Les 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
bundle async_bundle avec n'importe quelle combinaison de tenseurs classés de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 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 4/8/16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/ Entier signé quantifié uniforme de 32 bits ou entier non signé quantifié uniforme de 4/8/16/32 bits ou quantifié uniforme de 4/8/16/32 bits par axe entier signé ou quantifié uniforme de 4/8/16/32 bits par valeurs entières non signées de l'axe ou valeurs de jeton

Résultats:

Résultat Description
"anonyme" async_bundle avec n'importe quelle combinaison de tenseurs classés de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 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 4/8/16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou 64 bits ou 4/8/16/ Entier signé quantifié uniforme de 32 bits ou entier non signé quantifié uniforme de 4/8/16/32 bits ou quantifié uniforme de 4/8/16/32 bits par axe entier signé ou quantifié uniforme de 4/8/16/32 bits par valeurs entières non signées de l'axe 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou type complexe avec des éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 4/8/16/32 bits ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits
rhs tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou type complexe avec des éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 4/8/16/32 bits ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits

Résultats:

Résultat Description
result tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou flottant 16 bits ou flottant 32 bits ou flottant 64 bits ou type bfloat16 ou type complexe avec des éléments flottants 32 bits ou 64 bits ou Entier signé quantifié uniforme 4/8/16/32 bits ou valeurs entières non signées quantifiées uniformes 4/8/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{}

Les attributs:

Attribut Type MLIR Description
epsilon ::mlir::FloatAttr Attribut float 32 bits
feature_index ::mlir::IntegerAttr Attribut entier sans signe de 64 bits

Opérandes :

Opérande Description
operand Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
scale Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
mean Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
variance Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
grad_output Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16

Résultats:

Résultat Description
grad_operand Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
grad_scale Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
grad_offset Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de 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{}

Les attributs:

Attribut Type MLIR Description
epsilon ::mlir::FloatAttr Attribut float 32 bits
feature_index ::mlir::IntegerAttr Attribut entier sans signe de 64 bits

Opérandes :

Opérande Description
operand Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
scale Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
offset Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
mean Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
variance Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de 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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float 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{}

Les attributs:

Attribut Type MLIR Description
epsilon ::mlir::FloatAttr Attribut float 32 bits
feature_index ::mlir::IntegerAttr Attribut entier sans signe de 64 bits

Opérandes :

Opérande Description
operand Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
scale Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
offset Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16

Résultats:

Résultat Description
output Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
batch_mean Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16
batch_var Tenseur 1D de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou de valeurs de type float 16 bits ou 32 bits float ou 64 bits float ou bfloat16

mhlo.bitcast (mhlo :: BitcastOp)

Opération de diffusion de bits

Syntaxe:

operation ::= `mhlo.bitcast` 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 modifie la forme de l'entrée de manière à ce que la disposition physique des éléments reste inchangée.

Cette opération nécessite des informations de mise en page pour donner un sens à « la disposition physique des éléments », et la prise en charge de la mise en page dans MHLO est actuellement en cours de développement.

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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits par axe

Résultats:

Résultat Description
"anonyme" Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits 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 dans lequel 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/spec.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 f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 Entier sans signe /16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou entier non signé quantifié uniforme 4/8/16/32 bits ou quantifié uniforme 4/8/16/32 bits par axe entier signé ou valeurs entières non signées quantifiées uniformes 4/8/16/32 bits par axe

Résultats:

Résultat Description
"anonyme" Tenseur classé de type f8E4M3B11FNUZ ou de type f8E4M3FN ou de type f8E4M3FNUZ ou de type f8E5M2 ou de type f8E5M2FNUZ ou float 16 bits ou float 32 bits ou float 64 bits ou type bfloat16 ou pred (AKA booléen ou entier 1 bit) ou 4/8 /Entier sans signe 16/32/64 bits ou entier non signé 4/8/16/32/64 bits ou type complexe avec éléments flottants 32 bits ou flottants 64 bits ou quantification uniforme 4/8/16/32 bits entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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

Les 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 Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
"anonyme" Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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

Les 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 Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
"anonyme" Tenseur de forme statique de type F8E4M3B11fnuz ou type F8E4M3fn ou F8E4M3Fnuz ou Type F8E5M2 ou Float F8E5M2FNUZ OU Float 16 bits OU Float 32 bits ou Float 64 bits ou 4/4/4/4/4 8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 4/8/16/32 bits uniformes entier signé quantifié ou 4/8/16/32 bits uniformes entières non signées quantifiées ou 4/8/16/32 bits uniformes quantisées par axe entier signé ou valeurs entières 4/8/16/32 bits par axe quantisé par axe non signé non signé

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" Variadique du tenseur classé de type F8E4M3B11fnuz ou type F8E4M3FN ou F8E4M3FNUZ OU TYPE F8E5M2 ou F8E5M2FNUZ OU Float 16 bits ou Float 32 bits ou Float 64 bits ou integer BFLOAT16 / 8/16/32/64 bits entier sans signe ou 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 4/8/16/32 bits Uniforme entier signé quantifié ou 4/8/16/32 bits Valeurs entières uniques uniques unifiées ou tenseur classé de 4/8/16/32 bits Uniforme quantisé par axe signé entier ou 4/8/16/32 bits 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 , Elementwise , SameOperandsAndResultShape

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

Effets: MemoryEffects::Effect{}

Opérandes:

Opérande Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32 bits uniformes entier signé quantifié ou 4/8/16/32 bits

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32 bits uniformes entier signé quantifié ou 4/8/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 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 , Elementwise , SameOperandsAndResultShape

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

Effets: MemoryEffects::Effect{}

Opérandes:

Opérande Description
operand Tenseur classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou type F8E4M3fnuz ou type F8E5M2 ou Float F8E5M2FNUZ OU FLOT 16 bits ou flotteur 32 bits ou Valeurs de type BFloat16 à 64 bits ou bfloat16

Résultats:

Résultat Description
result Tenseur classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou type F8E4M3fnuz ou type F8E5M2 ou Float F8E5M2FNUZ OU FLOT 16 bits ou flotteur 32 bits ou Valeurs de type BFloat16 à 64 bits ou bfloat16

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

Les attributs:

Attribut Type mlir Description
lower :: mlir :: boolattr attribut bool

Opérandes:

Opérande Description
a Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3FN ou F8E4M3Fnuz ou F8E5M2 ou F8E5M2FNUZ ou Type ou type de complexe 32 bits ou 64 bits Float ou BFLOAT16

Résultats:

Résultat Description
"anonyme" Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3FN ou F8E4M3Fnuz ou F8E5M2 ou F8E5M2FNUZ ou Type ou type de complexe 32 bits ou 64 bits Float ou BFLOAT16

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 , HLO_BroadcastingElementwise , InferTensorType , SameOperandsAndResultElementType

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

Effets: MemoryEffects::Effect{}

Opérandes:

Opérande Description
min Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits
operand Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits
max Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
result Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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

Les 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 Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
"anonyme" Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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

Les 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 Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
"anonyme" Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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 , Elementwise , InferTensorType , SameOperandsAndResultShape , SameOperandsElementType

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

Effets: MemoryEffects::Effect{}

Les 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 Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits
rhs Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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

Les 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 Variadique du tenseur classé de type F8E4M3B11fnuz ou type F8E4M3FN ou F8E4M3FNUZ OU TYPE F8E5M2 ou F8E5M2FNUZ OU Float 16 bits ou Float 32 bits ou Float 64 bits ou integer BFLOAT16 / 8/16/32/64 bits entier sans signe ou 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 4/8/16/32 bits entier signé quantifié uniforme ou 4/8/16/32 bits entier uniforme unifié ou signé ou 4/8/16/32 bits UNIFORMATION Per-Axe entier signé entier non signé 4/8/16/16 / 32 bits Valeurs ou jeton ou tuple imbriqué avec n'importe quelle combinaison de tenseur classé de type F8E4M3B11Fnuz ou de type F8E4M3FN ou F8E4M3FNUZ ou F8E5M2 ou Type F8E5M2FNUZ OU TYPE 16-BIT OU FLOOT 32-BIT OU FLOOT 64 BIT OU BFLOAT16 ou entier 1 bits) ou 4/8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 4 / 8/16/32 bits uniformes entier signé quantifié ou 4/8/16/32 bits Uniform Quantized Néantisé Valeurs entières non signées ou tenseur classé de 4/8/8/16/32 bits Quantisé par axe entier signé ou 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" Variadique du tenseur classé de type F8E4M3B11fnuz ou type F8E4M3FN ou F8E4M3FNUZ OU TYPE F8E5M2 ou F8E5M2FNUZ OU Float 16 bits ou Float 32 bits ou Float 64 bits ou integer BFLOAT16 / 8/16/32/64 bits entier sans signe ou 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 4/8/16/32 bits entier signé quantifié uniforme ou 4/8/16/32 bits entier uniforme unifié ou signé ou 4/8/16/32 bits UNIFORMATION Per-Axe entier signé entier non signé 4/8/16/16 / 32 bits Valeurs ou jeton ou tuple imbriqué avec n'importe quelle combinaison de tenseur classé de type F8E4M3B11Fnuz ou de type F8E4M3FN ou F8E4M3FNUZ ou F8E5M2 ou Type F8E5M2FNUZ OU TYPE 16-BIT OU FLOOT 32-BIT OU FLOOT 64 BIT OU BFLOAT16 ou entier 1 bits) ou 4/8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 4 / 8/16/32 bits uniformes entier signé quantifié ou 4/8/16/32 bits Uniform Quantized Néantisé Valeurs entières non signées ou tenseur classé de 4/8/8/16/32 bits Quantisé par axe entier signé ou 4/8 / 16/32 bits uniformes quantifiés par axe Valeurs entières non signées ou valeurs de jeton

mhlo.compute_reshape_shape (mhlo :: ComputerShapeshapeop)

Opération de contrebande informatique

Syntaxe:

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

Cette opération est un travail en cours, il n'est donc pas encore inclus dans la spécification: https://github.com/openxla/stablehlo/issues/8

De manière informelle, cette opération calcule un Output_Shape pour DynamicReshapeop à partir du nombre de num_elements d'éléments dans un opérande de DynamicReshapeop et de la forme dynamic_shape fournies à TF Reshape: https://www.tensorflow.org/api_docs/python/tf/reshape

Par exemple, pour num_elements = 12 et dynamic_shape = [2, -1] , le result va être [2, 6] . Si les opérandes ne sont pas valides (par exemple, si les dimensions ne divisent pas uniformément le nombre d'éléments, ou s'il y a plusieurs valeurs -1 dans les dimensions), cela conduit à un comportement non défini.

Exemple:

%result = mhlo.compute_reshape_shape %num_elements, %dynamic_shape
       : (index, tensor<2xi32>) -> tensor<2xi32>

TRAITS: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effets: MemoryEffects::Effect{}

Opérandes:

Opérande Description
num_elements indice
dynamic_shape Tenseur 1d de valeurs entières ou d'index

Résultats:

Résultat Description
result Tenseur 1d de valeurs entières ou d'index

mhlo.concatenate (mhlo :: concatenateop)

Opération de concaténate

COMPATENNES UN NOMBRE VARIADIQUE DE TENSEURS DANS inputs dimension DIMENDENCE 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{}

Les attributs:

Attribut Type mlir Description
dimension :: Mlir :: Integerattr Attribut entier sans signe 64 bits

Opérandes:

Opérande Description
val Variadique du tenseur classé de type F8E4M3B11fnuz ou type F8E4M3FN ou F8E4M3FNUZ OU TYPE F8E5M2 ou F8E5M2FNUZ OU Float 16 bits ou Float 32 bits ou Float 64 bits ou integer BFLOAT16 / 8/16/32/64 bits entier sans signe ou 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 4/8/16/32 bits entier signé quantifié uniforme ou 4/8/16/32 bits entier uniforme unifié ou signé ou 4/8/16/32 bits UNIFORMATION Per-Axe entier signé entier non signé 4/8/16/16 / 32 bits valeurs

Résultats:

Résultat Description
"anonyme" Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

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

Les attributs:

Attribut Type mlir Description
value :: mlir :: elementSattr Attribut vectoriel / tensor constant

Résultats:

Résultat Description
output Tenseur de forme statique de type F8E4M3B11fnuz ou type F8E4M3fn ou F8E4M3Fnuz ou Type F8E5M2 ou Float F8E5M2FNUZ OU Float 16 bits OU Float 32 bits ou Float 64 bits ou 4/4/4/4/4 8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 4/8/16/32 bits uniformes entier signé quantifié ou 4/8/16/32 bits uniformes entières non signées quantifiées ou 4/8/16/32 bits uniformes quantisées par axe entier signé ou valeurs entières 4/8/16/32 bits par axe quantisé par axe non signé non signé

mhlo.convert (mhlo :: convertop)

Convertir

Syntaxe:

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

Effectue une conversion d'élément d'un type d'élément à un autre sur le tenseur operand et produit un tenseur result .

Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.md#convert

Exemple:

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

TRAITS: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effets: MemoryEffects::Effect{}

Opérandes:

Opérande Description
operand Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
result Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

mhlo.convolution (Mhlo :: Convolutionop)

Opération de convolution

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)

Calcule des produits DOT entre Windows de lhs et des tranches de rhs et produit result .

Voir: https://github.com/openxla/stablehlo/blob/main/docs/spe.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)

Effets: MemoryEffects::Effect{}

Les attributs:

Attribut Type mlir Description
window_strides :: mlir :: dense intelementsAtTr Attribut d'éléments entiers sans signe 64 bits
padding :: mlir :: dense intelementsAtTr Attribut d'éléments entiers sans signe 64 bits
lhs_dilation :: mlir :: dense intelementsAtTr Attribut d'éléments entiers sans signe 64 bits
rhs_dilation :: mlir :: dense intelementsAtTr Attribut d'éléments entiers sans signe 64 bits
window_reversal :: mlir :: denseelementsattr Vector booléen constant / Tensor Attribut
dimension_numbers :: Mlir :: Mhlo :: ConvdimensionNumbersattr Structure des informations sur la dimension pour Conv Op
feature_group_count :: Mlir :: Integerattr Attribut entier sans signe 64 bits
batch_group_count :: Mlir :: Integerattr Attribut entier sans signe 64 bits
precision_config :: mlir :: arrayattr Attribut de configuration de précision

Opérandes:

Opérande Description
lhs Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits
rhs Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

Résultats:

Résultat Description
"anonyme" Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits entier uniforme non signé ou 4/8/16/16/32 bits quantifié par axe entier signé ou 4/8/16/32 bits

mhlo.copy (mhlo :: copyop)

Opération de copie

Syntaxe:

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

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 une copie de operand . Selon les métadonnées attachées à l'opération, il peut se comporter très différemment d'un no-opér.

Exemple:

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

TRAITS: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

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

Effets: MemoryEffects::Effect{}

Les attributs:

Attribut Type mlir Description
cross_program_prefetch_index :: Mlir :: Integerattr Attribut entier sans signe 32 bits

Opérandes:

Opérande Description
operand Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits uniformes entier non signé quantifié ou 4/8/16/32 bits uniformes quantisées par axe entier signé ou 4/8/16/32 bits UNIFICAL MANDIFED Perte Token ou tuple imbriqué avec n'importe quelle combinaison de tenseur classé de type F8E4M3B11Fnuz ou de type F8E4M3FN ou F8E4M3Fnuz ou Type F8E5M2 OU F8E5M2FNUZ OU FLOTOY -bit entier) ou 4/8/16/32/64 bits entier sans signe ou 4/8/16/32/64 bits entier non signé ou type complexe avec flotteur 32 bits ou éléments flottants 64 bits ou 4/8 / 16/32 bits uniformes entières signées quantifiées ou 4/8/16/32 bits Valeurs entières non signées quantifiées ou tenseur classé de 4/8/16/16 / 32 bits Quantisé par axe entier signé ou 4/8/16 / 32 bits Uniform Qualzed Per Axe Valeurs entières non signées ou valeurs de jeton

Résultats:

Résultat Description
result Tensor classé de type F8E4M3B11Fnuz ou type F8E4M3Fn ou F8E4M3Fnuz ou Type F8E5M2 ou F8e5m2fnuz ou flotteur 16 bits ou float 32 bits ou intime 64 bits ou 4/8 bfloat16 / 16/32/64 bits entier sans signe ou 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 4/8/16/32 bits quantifiés uniformes quantifiés entier signé ou 4/8/16/32 bits uniformes entier non signé quantifié ou 4/8/16/32 bits uniformes quantisées par axe entier signé ou 4/8/16/32 bits UNIFICAL MANDIFED Perte Token ou tuple imbriqué avec n'importe quelle combinaison de tenseur classé de type F8E4M3B11Fnuz ou de type F8E4M3FN ou F8E4M3Fnuz ou Type F8E5M2 OU F8E5M2FNUZ OU FLOTOY -bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8 /16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 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{}

Les attributs:

Attribut MLIR Type Description
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.cstr_reshapable (mhlo::CstrReshapableOp)

CstrReshapable operation

Syntaxe:

operation ::= `mhlo.cstr_reshapable` 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 creates a witness on the constraint that ComputeReshapeShape would succeed with the provided operands.

Exemple:

%result = mhlo.cstr_reshapable %num_elements, %dynamic_shape
       : (index, tensor<3xi32>) -> !shape.witness

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand Description
num_elements indice
dynamic_shape 1D tensor of integer or index values

Résultats:

Résultat Description
result

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

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Operand Description
lhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

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

Operands:

Operand Description
lhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_broadcast_in_dim (mhlo::DynamicBroadcastInDimOp)

DynamicBroadcastInDim 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 BroadcastInDimOp except that the result shape is specified dynamically via output_dimensions : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#broadcast_in_dim

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.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
output_dimensions 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

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
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Operand Description
lhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
d_padding ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
slice_sizes ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_iota (mhlo::DynamicIotaOp)

DynamicIota 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 IotaOp except that the result shape is specified dynamically via output_shape : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#iota

Exemple:

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Les attributs:

Attribut MLIR Type Description
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

Operand Description
output_shape 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
edge_padding_low 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
edge_padding_high 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
interior_padding 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 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 ReshapeOp except that the result shape is specified dynamically via output_shape : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reshape

Exemple:

%0 = mhlo.dynamic_reshape %arg0, %shape : (tensor<?xf32>, tensor<2xindex>) -> tensor<?x?xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand Description
operand tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
output_shape 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
update ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
einsum_config ::mlir::StringAttr string attribute

Operands:

Operand Description
lhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type 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.

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or nested tuple with any combination of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 = #mhlo.gather<
    offset_dims = [2, 3],
    collapsed_slice_dims = [0],
    start_index_map = [0, 2],
    index_vector_dim = 2>,
  slice_sizes = dense<[0, 2, 2]> : tensor<3xi64>,
  indices_are_sorted = false
} : (tensor<3x4x2xi32>, tensor<2x3x2xi64>) -> tensor<2x3x2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

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

Effects: MemoryEffects::Effect{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
dimension ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
index ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

Operand Description
operand nested tuple with any combination of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.if (mhlo::IfOp)

If operation

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

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

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

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface

Operands:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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)

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute

Résultats:

Résultat Description
output statically shaped tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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

Les attributs:

Attribut MLIR Type Description
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Description
inputs variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 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 4/8/16/32/64-bit signless integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

Attribut MLIR Type Description
outfeed_config ::mlir::StringAttr string attribute

Operands:

Operand Description
inputs variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs
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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
limit_indices 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
strides 1D tensor of index or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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)

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

Attribut MLIR Type Description
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Description
inputs variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs
init_values variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

Résultats:

Résultat Description
"anonyme" variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

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

Les attributs:

Attribut MLIR Type Description
exponent_bits ::mlir::IntegerAttr 32-bit signless integer attribute
mantissa_bits ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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>

Les attributs:

Attribut MLIR Type Description
scatter_dimension ::mlir::IntegerAttr 64-bit signless integer attribute
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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs
init_values variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

Résultats:

Résultat Description
"anonyme" variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" statically shaped tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
shape 1D tensor of index or 4/8/16/32/64-bit signless integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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{}

Les attributs:

Attribut MLIR Type Description
rng_algorithm ::mlir::mhlo::RngAlgorithmAttr XLA PRNG algorithm to be used.

Operands:

Operand Description
initial_state ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
output statically shaped tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 = [2,3],
    inserted_window_dims = [0],
    scatter_dims_to_operand_dims = [1, 0],
    index_vector_dim = 2>,
  indices_are_sorted = false,
  unique_indices = false
} : (tensor<3x4x2xi32>, tensor<2x3x2xi64>, tensor<2x3x2x2xi32>) -> tensor<3x4x2xi32>

Traits: RecursiveMemoryEffects , SameVariadicOperandSize

Interfaces: InferTypeOpInterface

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs
scatter_indices ranked tensor of integer or index values
updates variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

Résultats:

Résultat Description
"anonyme" variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

mhlo.select (mhlo::SelectOp)

Sélectionnez l'opération

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
on_false ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
source ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
init_value ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs
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{}

Les attributs:

Attribut MLIR Type Description
dimension ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 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 4/8/16/32/64-bit signless integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

Résultats:

Résultat Description
"anonyme" variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

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

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
meta variadic of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4 /8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer valeurs

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
random ranked tensor of 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 4/8/16/32-bit uniform quantized signed integer or 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

Les attributs:

Attribut MLIR Type Description
k ::mlir::IntegerAttr 64-bit signless integer attribute
largest ::mlir::BoolAttr bool attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
values ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
indices ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values
index ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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>

Les attributs:

Attribut MLIR Type Description
tag ::mlir::StringAttr string attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

Attribut MLIR Type Description
permutation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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{}

Les attributs:

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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.unary_einsum (mhlo::UnaryEinsumOp)

UnaryEinsum 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.unary_einsum"(%operand) {
  einsum_config = "ab->a"
} : (tensor<4x16xf32>) -> tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Les attributs:

Attribut MLIR Type Description
einsum_config ::mlir::StringAttr string attribute

Operands:

Operand Description
operand ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
"anonyme" ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer 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 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Résultats:

Résultat Description
result ranked tensor of 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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 f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 4/8/16/32-bit uniform quantized per axis signed integer or 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

Les attributs:

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 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer values

Résultats:

Résultat Description
result ranked tensor of f8E4M3B11FNUZ type or f8E4M3FN type or f8E4M3FNUZ type or f8E5M2 type or f8E5M2FNUZ type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 4/8/16/32/64-bit signless integer or 4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 4/8/16/32-bit uniform quantized signed integer or 4/8/16/32-bit uniform quantized unsigned integer or 4/8/16/32-bit uniform quantized per axis signed integer or 4/8/16/32-bit uniform quantized per axis unsigned integer values

Les 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 ) #### Parameters:
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 ) #### Parameters:
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 ) #### Parameters:
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 ) #### Parameters:
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 ) #### Parameters:
Paramètre C++ type Description
valeur ::mlir::mhlo::DomainKind an enum of type DomainKind

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

Les types

AsyncBundleType

Opaque collection of other types

Syntaxe:

!mhlo.async_bundle<
  ::llvm::ArrayRef<Type>   # types
>

Paramètres:

Paramètre C++ type Description
les types ::llvm::ArrayRef<Type>