Spécification StableHLO

StableHLO est un ensemble d'opérations pour les opérations de haut niveau (HLO) dans les modèles de ML (ML). StableHLO fonctionne comme une couche de portabilité entre différents Frameworks et compilateurs de ML: frameworks de ML qui produisent des programmes StableHLO sont compatibles avec les compilateurs de ML qui consomment des programmes StableHLO.

Notre objectif est de simplifier et d'accélérer le développement du ML en créant l'interopérabilité entre différents frameworks de ML (tels que TensorFlow, JAX et PyTorch) et les compilateurs de ML (tels que XLA et IREE). Dans cette optique, fournit une spécification pour le langage de programmation StableHLO.

Cette spécification contient trois sections principales. Tout d'abord, le La section Programmes décrit la structure des programmes StableHLO. qui se composent de fonctions StableHLO qui sont elles-mêmes des opérations StableHLO. Au sein de cette structure, la section Ops spécifie la sémantique opérations individuelles. La section Exécution fournit une sémantique pour toutes ces opérations s'exécutent ensemble dans un programme. Enfin, la fonction La section Notation présente la notation utilisée tout au long de la spécifique.

Pour afficher la spécification d'une version précédente de StableHLO, ouvrez le dépôt à l'adresse version taguée qui vous intéresse. (par exemple, la spécification StableHLO v0.19.0). Pour afficher les modifications survenues lors de chaque chargement de version mineure de StableHLO, consultez le journal des versions dans VhloDialect.td.

Programmes

Program ::= {Func}

Les programmes StableHLO consistent en un nombre arbitraire de fonctions StableHLO. Vous trouverez ci-dessous un exemple de programme avec une fonction @main comportant trois entrées (%image, %weights et %bias) et un résultat. Corps de la fonction comporte 6 opérations.

func.func @main(
  %image: tensor<28x28xf32>,
  %weights: tensor<784x10xf32>,
  %bias: tensor<1x10xf32>
) -> tensor<1x10xf32> {
  %0 = "stablehlo.reshape"(%image) : (tensor<28x28xf32>) -> tensor<1x784xf32>
  %1 = "stablehlo.dot"(%0, %weights) : (tensor<1x784xf32>, tensor<784x10xf32>) -> tensor<1x10xf32>
  %2 = "stablehlo.add"(%1, %bias) : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32>
  %3 = "stablehlo.constant"() {value = dense<0.0> : tensor<1x10xf32>} : () -> tensor<1x10xf32>
  %4 = "stablehlo.maximum"(%2, %3) : (tensor<1x10xf32>, tensor<1x10xf32>) -> tensor<1x10xf32>
  "func.return"(%4): (tensor<1x10xf32>) -> ()
}

Fonctions

Func        ::= 'func' '.' 'func' FuncId FuncInputs FuncOutputs '{' FuncBody '}'
FuncInputs  ::= '(' [FuncInput {',' FuncInput}] `)`
FuncInput   ::= ValueId ':' ValueType
FuncOutputs ::= ['->' FuncOutput, {',' FuncOutput}]
FuncOutput  ::= ValueType
FuncBody    ::= {Op}

Les fonctions StableHLO (également appelées fonctions nommées) ont un identifiant, des entrées/sorties et un corps. À l'avenir, nous prévoyons de Intégrer des métadonnées supplémentaires pour les fonctions afin d'améliorer la compatibilité avec HLO (#425, #626, n° 740, n° 744).

Identifiants

FuncId  ::= '@' letter {letter | digit}
ValueId ::= '%' digit {digit}
          | '%' letter {letter | digit}
letter  ::= 'a' | ... | 'z' | 'A' | ... | 'Z' | '_'
digit   ::= '0' | ... | '9'

Les identifiants StableHLO sont semblables aux identifiants dans de nombreux programmes avec deux particularités: 1) tous les identifiants ont des sigils distinguer différents types d'identifiants, 2) les identifiants de valeur peuvent être entièrement numérique pour simplifier la génération de programmes StableHLO.

Types

Type         ::= ValueType | NonValueType
ValueType    ::= TensorType | QuantizedTensorType | TokenType | TupleType
NonValueType ::= TensorElementType | QuantizedTensorElementType | FunctionType | StringType

Les types StableHLO sont classés en types de valeurs (également appelés types de première classe), qui représentent des valeurs StableHLO et des types autres que des valeurs qui décrivent d'autres éléments du programme. Les types StableHLO sont similaires aux types dans de nombreux langages de programmation, la principale particularité étant domaine spécifique qui donne des résultats inhabituels (par exemple, des types scalaires ne sont pas des types de valeur).

TensorType ::= 'tensor' '<' Shape TensorElementType '>'
Shape ::= {DimensionSize 'x'}
DimensionSize ::= digit {digit} | '?'

Les types de Tensor représentent des Tensors, c'est-à-dire des tableaux multidimensionnels. Ils ont un forme et un type d'élément, où une forme représente des valeurs non négatives ou tailles de dimensions inconnues dans l'ordre croissant des valeurs dimensions (également appelées axes) numérotées de 0 à R-1. La Le nombre de dimensions R est appelé classement. Par exemple, tensor<2x3xf32> est un type de Tensor avec la forme 2x3 et le type d'élément f32. Il a deux dimensions (ou, en d'autres termes, deux axes) - 0e dimension et 1re dimension - dont les tailles sont 2 et 3. Son classement est de 2.

Les formes peuvent être partiellement ou totalement inconnues (dynamiques). Ex. : tensor<?x2xf64> est en partie inconnue, tandis que tensor<?x?xf64> l'est complètement. Dynamique les dimensions sont représentées à l'aide d'un ?. Les formes ne peuvent pas être désclassées.

À l'avenir, nous prévoyons d'étendre les types de Tensor au-delà les tailles de dimension et les types d'éléments, par exemple, pour inclure des mises en page (#629) et la parcimonie (#1078)

QuantizedTensorType ::= 'tensor' '<' Shape QuantizedTensorElementType '>'
QuantizedTensorElementType ::= '!quant.uniform' '<'
                  QuantizationStorageType
                  ['<' QuantizationStorageMin ':' QuantizationStorageMax '>']
                  ':' QuantizationExpressedType
                  [':' QuantizationDimension]
                  ',' QuantizationParameters '>'
QuantizationStorageType ::= IntegerType
QuantizationStorageMin ::= IntegerConstant
QuantizationStorageMax ::= IntegerConstant
QuantizationExpressedType ::= FloatType
QuantizationDimension ::= IntegerConstant
QuantizationParameters ::= QuantizationParameter
                         | '{' QuantizationParameter {',' QuantizationParameter} '}'
QuantizationParameter ::= QuantizationScale ':' QuantizationZeroPoint
QuantizationScale ::= FloatConstant
QuantizationZeroPoint ::= IntegerConstant

Les types d'éléments quantifiés représentent des valeurs entières d'un type de stockage dans la plage comprise entre storage_min et storage_max (inclus) correspondant à valeurs à virgule flottante d'un type exprimé. Pour une valeur entière donnée i, la valeur à virgule flottante correspondante f peut être calculée comme suit : f = (i - zero_point) * scale, où scale et zero_point sont appelés paramètres de quantification. Les champs storage_min et storage_max sont facultatifs dans la grammaire, mais leur valeur par défaut est min_value(storage_type) et max_value(storage_type) respectivement. Les types d'éléments quantifiés ont la les contraintes suivantes:

• (C1) type(storage_min) = storage_type.
• (C2) type(storage_max) = storage_type.
• (C3) min_value(storage_type) <= storage_min < storage_max <= max_value(storage_type).
• (C4) type(scales...) = expressed_type.
• (C5) 0 < scales.
• (C6) is_finite(scales...).
• (C7) storage_min <= zero_points <= storage_max.
• (C8) type(zero_points...) = storage_type.
• (C9) size(scales) = size(zero_points).
• (C10) Si is_empty(quantization_dimension), alors size(scales) = 1.
• (C11) 0 <= quantization_dimension.

Pour le moment, QuantizationScale est une constante à virgule flottante, mais il existe un fort intérêt pour les échelles basées sur des entiers, représentées par des multiplicateurs et changements de direction. Nous prévoyons de l'examiner prochainement. (#1404)

Une discussion est en cours sur la sémantique de QuantizationZeroPoint, y compris le type, les valeurs et s'il ne peut y avoir qu'un seul ou potentiellement plusieurs points zéro dans un type de Tensor quantifié. D'après les résultats de cette discussion, la spécification autour de zéro point peut changer ultérieurement (#1405).

Une autre discussion en cours concerne la sémantique de QuantizationStorageMin. et QuantizationStorageMax pour déterminer si des contraintes doivent être sur ces valeurs et sur les valeurs des Tensors quantifiés (#1406)

Enfin, nous prévoyons de représenter les échelles inconnues et les valeurs zéro de la même manière que nous prévoyons d'explorer la représentation des (#1407)

Les types de Tensors quantifiés représentent des Tensors avec des éléments quantifiés. Ces Les Tensors sont exactement les mêmes que les Tensors standards, si ce n'est que leurs éléments utilisent des types d'éléments quantifiés au lieu de types d'éléments standards.

Dans les tenseurs quantifiés, la quantification peut être par Tensor, ce qui signifie que une valeur scale et une zero_point pour l'intégralité du Tensor ou par axe, c'est-à-dire qu'avoir plusieurs scales et zero_points, une paire par tranche de une dimension particulière quantization_dimension. Plus formellement, dans un Tensor t avec la quantification par axe, il existe dim(t, quantization_dimension) tranches pour quantization_dimension: t[:, ..., 0, ..., :], t[:, ..., 1, ..., :], etc. Tous les éléments de la ie tranche utilisent scales[i] et zero_points[i] comme leurs paramètres de quantification. Les types de Tensor quantifiés ont les caractéristiques suivantes : contraintes:

• Pour la quantification par Tensor: <ph type="x-smartling-placeholder">
  </ph>
  • Aucune contrainte supplémentaire.
• Pour la quantification par axe: <ph type="x-smartling-placeholder">
  </ph>
  • (C12) quantization_dimension < rank(self).
  • (C13) dim(self, quantization_dimension) = size(scales).

TokenType ::= 'token'

Les types de jetons représentent les jetons, c'est-à-dire les valeurs opaques produites et consommées par certaines opérations. Les jetons sont utilisés pour imposer un ordre d'exécution aux opérations comme décrit dans la section Exécution.

TupleType ::= 'tuple' '<' TupleElementTypes '>'
TupleElementTypes ::= [ValueType {',' ValueType}]

Les types de tuples représentent des tuples, c'est-à-dire des listes hétérogènes. Les Tuples sont un héritage fonctionnalité qui n'existe que pour la compatibilité avec HLO. Dans HLO, les tuples sont utilisée pour représenter des entrées et sorties variables. Dans StableHLO, les entrées variables et les sorties sont prises en charge de manière native, et la seule utilisation de tuples dans StableHLO consiste à représenter de manière exhaustive l'ABI HLO T, tuple<T> et tuple<tuple<T>> peut être sensiblement différente en fonction d'une la mise en œuvre. Nous prévoyons d'apporter des modifications à l'ABI HLO à l'avenir. ce qui peut nous permettre de supprimer les types de tuples de StableHLO (#598)

TensorElementType ::= BooleanType | IntegerType | FloatType | ComplexType
BooleanType ::= 'i1'
IntegerType ::= SignedIntegerType | UnsignedIntegerType
SignedIntegerType ::= 'si2' | 'si4' | 'si8' | 'si16' | 'si32' | 'si64'
UnsignedIntegerType ::= 'ui2' | 'ui4' | 'ui8' | 'ui16' | 'ui32' | 'ui64'
FloatType ::= 'f8E4M3FN' | 'f8E5M2' | 'f8E4M3FNUZ' | 'f8E5M2FNUZ'
            | 'f8E4M3B11FNUZ' | 'bf16' | 'f16' | 'f32' | 'f64'
TensorFloat32 ::= 'tf32'
ComplexType ::= 'complex' '<' ComplexElementType '>'
ComplexElementType ::= 'f32' | 'f64'

Les types d'éléments représentent les éléments des types de Tensor. Contrairement à de nombreuses fonctions langues, ces types ne sont pas de première classe dans StableHLO. Cela signifie que Les programmes StableHLO ne peuvent pas représenter directement des valeurs de ces types (par conséquent, représenter des valeurs scalaires de type T avec un Tensor à 0 dimensions est idiomatique. de type tensor<T>).

• Le type booléen représente les valeurs booléennes true et false.
• Les types d'entiers peuvent être signés (si) ou non signés (ui), et ont l'une des largeurs de bits acceptées (2, 4, 8, 16, 32 ou 64) ; Les types siN signés représentent des valeurs entières comprises entre -2^(N-1) et 2^(N-1)-1 inclus, et les types uiN non signés représentent des valeurs entières comprises entre 0 et 2^N-1 inclus.
• Il existe plusieurs types à virgule flottante: <ph type="x-smartling-placeholder">
  </ph>
  • les types f8E4M3FN et f8E5M2 correspondant respectivement aux Les encodages E4M3 et E5M2 du format FP8 décrits à la section Formats FP8 pour le deep learning.
  • Types f8E4M3FNUZ et f8E5M2FNUZ correspondant à E4M3 et E5M2 les encodages des formats FP8 décrits dans Formats numériques 8 bits pour les réseaux de neurones profonds
  • Type de f8E4M3B11FNUZ correspondant à l'encodage E4M3 des formats FP8 décrits dans Entraînement et inférence hybrides à virgule flottante 8 bits (HFP8) pour les réseaux de neurones profonds
  • Type bf16 correspondant au format bfloat16 décrit dans BFloat16: Le secret des hautes performances sur les Cloud TPU
  • les types f16, f32 et f64 correspondant respectivement à binary16 ("demi-précision"), binary32 ("précision simple") et les formats binary64 ("double précision") décrits dans les la norme IEEE 754.
  • Le type de tf32 correspond au format TensorFloat32. et offre une compatibilité limitée avec StableHLO.
• Les types complexes représentent des valeurs complexes ayant une partie réelle. et une partie imaginaire du même type d'élément. Complexe compatible les types sont complex<f32> (les deux parties sont de type f32) et complex<f64> (les deux parties sont de type f64).

FunctionType ::= '(' InputTypes ')' '->' '(' OutputTypes ')'
InputTypes ::= [ValueType {',' ValueType}]
OutputTypes ::= [ValueType {',' ValueType}]

Les types de fonctions représentent les fonctions nommées et anonymes. Ils ont les types d'entrée (liste des types à gauche de ->) et les types de sortie (la liste des types sur la droite de ->). Dans de nombreux langages de programmation, langages, les types de fonctions sont de première classe, mais pas dans StableHLO.

StringType ::= 'string'

Le type de chaîne représente des séquences d'octets. Contrairement à de nombreuses fonctions langues, le type de chaîne n'est pas la première classe dans StableHLO et n'est utilisé spécifier des métadonnées statiques pour les éléments du programme.

Opérations

Les opérations StableHLO (également appelées opérations) représentent un ensemble fermé. d'opérations de haut niveau dans des modèles de machine learning. Comme indiqué ci-dessus, La syntaxe StableHLO s'inspire fortement de MLIR, qui n'est pas nécessairement la méthode la plus alternative ergonomique, mais c'est sans doute la meilleure solution pour permettre à StableHLO de ce qui renforce l'interopérabilité entre les frameworks de ML et les compilateurs de ML.

Op            ::= [OpOutputs] OpName OpInputs ':' OpSignature
OpName        ::= '"' 'stablehlo' '.' OpMnemonic '"'
OpMnemonic    ::= 'abs' | 'add' | ...

Les opérations StableHLO (également appelées opérations) ont un nom, des entrées/sorties et une signature. Ce nom se compose du préfixe stablehlo. et Un mnémonique qui identifie de manière unique l'une des opérations prises en charge. Voir ci-dessous pour une liste complète de toutes les opérations prises en charge.

OpInputs        ::= OpInputValues OpInputFuncs OpInputAttrs
OpInputValues   ::= '(' [OpInputValue {',' OpInputValue}] ')'
OpInputValue    ::= ValueId
OpInputFuncs    ::= ['(' OpInputFunc {',' OpInputFunc} ')']
OpInputAttrs    ::= ['{' OpInputAttr {',' OpInputAttr} '}']
OpOutputs       ::= [OpOutput {',' OpOutput} '=']
OpOutput        ::= ValueId

Les opérations utilisent les entrées et produisent des sorties. Les entrées sont classées en les valeurs d'entrée (calculées lors de l'exécution), les fonctions d'entrée (fournies de manière statique, car dans StableHLO, les fonctions ne sont pas des valeurs de première classe) et d'entrée (également fournis de manière statique). Le type d'entrées et de sorties consommée et produite par une opération dépend de son mode mnémotechnique. Par exemple, add "op" consomme deux valeurs d'entrée et génère une valeur de sortie. En comparaison, les L'opération select_and_scatter consomme 3 valeurs d'entrée, 2 fonctions d'entrée et 3 attributs d'entrée.

OpInputFunc ::= '{' Unused FuncInputs ':' FuncBody '}'
Unused      ::= '^' digit {digit}
              | '^' letter {letter | digit}

Les fonctions d'entrée (également appelées fonctions anonymes) sont très semblables aux fonctions nommées, sauf que: 1) elles n'ont pas d'identifiant (donc le nom "anonyme"), 2) ils ne déclarent pas de types de sortie (les types de sortie sont inférée à partir de l'opération return dans la fonction).

La syntaxe des fonctions d'entrée inclut une partie actuellement inutilisée (voir les Unused ci-dessus), qui assure la compatibilité avec MLIR. Dans MLIR, il existe un concept plus général de « régions » qui peut comporter plusieurs "blocs" reliés entre eux par des jump ops. Ces blocs ont des ID qui correspondent à l'environnement de production Unused, afin de pouvoir les distinguer les uns des autres. StableHLO n'a pas de jump ops. La partie correspondante de la syntaxe MLIR est donc unused (mais est toujours là).

OpInputAttr      ::= OpInputAttrName '=' OpInputAttrValue
OpInputAttrName  ::= letter {letter | digit}
OpInputAttrValue ::= Constant

Les attributs d'entrée ont un nom et une valeur faisant partie des constantes. Elles constituent le principal moyen de spécifier des métadonnées statiques pour un programme éléments. Par exemple, l'opération concatenate utilise l'attribut dimension pour spécifier la dimension selon laquelle ses valeurs d'entrée sont concaténées. De même, L'opération slice utilise plusieurs attributs tels que start_indices et limit_indices. pour spécifier les limites utilisées pour segmenter la valeur d'entrée.

Actuellement, les programmes StableHLO contiennent parfois des attributs qui ne sont pas décrites dans ce document. À l'avenir, nous prévoyons de absorber ces attributs dans l'opset StableHLO ou les interdit dans les programmes StableHLO. En attendant, voici la liste de ces Attributs:

• layout (#629)
• mhlo.frontend_attributes (#628)
• mhlo.sharding (#619)
• output_operand_aliases (#740)
• Métadonnées de localisation (#594)

OpSignature ::= '(' [ValueType {',' ValueType}] ')' '->' '(' [ValueType {',' ValueType}] ')'

La signature d'opération comprend les types de toutes les valeurs d'entrée (la liste des types sur à gauche de ->) et les types de toutes les valeurs de sortie (liste des à droite de ->). À proprement parler, les types d'entrée sont redondants, et les types de sortie le sont presque toujours également (parce que pour la plupart des opérations StableHLO, les types de sortie peuvent être déduits des entrées). Néanmoins, l'opération la signature fait délibérément partie de la syntaxe StableHLO pour assurer la compatibilité avec MLIR.

Vous trouverez ci-dessous un exemple d'opération dont l'expression mnémotechnique est select_and_scatter. Il en consomme 3 valeurs d'entrée (%operand, %source et %init_value), deux fonctions d'entrée et trois attributs d'entrée (window_dimensions, window_strides et padding). Notez que la signature de l'opération n'inclut que les types de ses valeurs d'entrée. (mais pas les types de fonctions d'entrée et d'attributs qui sont fournis de façon intégrée).

%result = "stablehlo.select_and_scatter"(%operand, %source, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "stablehlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GE>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "stablehlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "stablehlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "stablehlo.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>

Constantes

Constant ::= BooleanConstant
           | IntegerConstant
           | FloatConstant
           | ComplexConstant
           | TensorConstant
           | QuantizedTensorConstant
           | StringConstant
           | EnumConstant

Les constantes StableHLO ont un littéral et un type qui représentent ensemble une valeur StableHLO. En général, le type fait partie de la syntaxe constante, sauf lorsqu'elle ne présente aucune ambiguïté (par exemple, une constante booléenne a sans ambiguïté le type i1, tandis qu'une constante entière peut avoir plusieurs types possibles).

BooleanConstant ::= BooleanLiteral
BooleanLiteral  ::= 'true' | 'false'

Les constantes booléennes représentent les valeurs booléennes true et false. Booléen sont de type i1.

IntegerConstant   ::= IntegerLiteral ':' IntegerType
IntegerLiteral    ::= ['-' | '+'] DecimalDigits
                    | ['-' | '+'] '0x' HexadecimalDigits
DecimalDigits     ::= decimalDigit {decimalDigit}
HexadecimalDigits ::= hexadecimalDigit {hexadecimalDigit}
decimalDigit      ::= '0' | ... | '9'
hexadecimalDigit  ::= decimalDigit | 'a' | ... | 'f' | 'A' | ... | 'F'

Les constantes entières représentent des valeurs entières via des chaînes utilisant des nombres décimaux ou en notation hexadécimale. Autres bases (par exemple, binaire ou octale, ne sont pas prises en charge. Les constantes entières présentent les contraintes suivantes:

• (C1) is_wellformed(integer_literal, integer_type).

FloatConstant  ::= FloatLiteral ':' FloatType
FloatLiteral   ::= SignPart IntegerPart FractionalPart ScientificPart
                 | '0x' [HexadecimalDigits]
SignPart       ::= ['-' | '+']
IntegerPart    ::= DecimalDigits
FractionalPart ::= ['.' [DecimalDigits]]
ScientificPart ::= [('e' | 'E') ['-' | '+'] DecimalDigits]

Les constantes à virgule flottante représentent des valeurs à virgule flottante via des chaînes qui utiliser la notation décimale ou scientifique. De plus, la notation hexadécimale peut être utilisé pour spécifier directement les bits sous-jacents au format à virgule flottante de le type correspondant. Les constantes à virgule flottante présentent les contraintes suivantes:

• (C1) Si vous utilisez la notation non hexadécimale, is_wellformed(float_literal, float_type)
• (C2) Si la notation hexadécimale est utilisée, size(hexadecimal_digits) = num_bits(float_type) / 4

ComplexConstant ::= ComplexLiteral ':' ComplexType
ComplexLiteral  ::= '(' RealPart ',' ImaginaryPart ')'
RealPart        ::= FloatLiteral
ImaginaryPart   ::= FloatLiteral

Les constantes complexes représentent des valeurs complexes à l'aide de listes d'une partie réelle (en premier) et une partie imaginaire (en deuxième). Par exemple : (1.0, 0.0) : complex<f32> représente 1.0 + 0.0i, et (0.0, 1.0) : complex<f32> représente 0.0 + 1.0i. L'ordre dans lequel ces sont ensuite stockées en mémoire et définies par l'implémentation. Constantes complexes présentent les contraintes suivantes:

• (C1) is_wellformed(real_part, complex_element_type(complex_type)).
• (C2) is_wellformed(imaginary_part, complex_element_type(complex_type)).

TensorConstant ::= TensorLiteral ':' TensorType
TensorLiteral  ::= 'dense' '<' (DenseLiteral | ElementLiteral) '>'
DenseLiteral   ::= DenseDimension | DenseElements
DenseDimension ::= '[' [DenseLiteral {',' DenseLiteral}] ']'
DenseElements  ::= [ElementLiteral {',' ElementLiteral}]
ElementLiteral ::= BooleanLiteral | IntegerLiteral | FloatLiteral | ComplexLiteral

Les constantes Tensor représentent les valeurs des Tensors au moyen de listes imbriquées spécifiées via Numpy. Exemple : dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> représente une valeur de Tensor avec le mappage suivant entre les index et les éléments: {0, 0} => 1, {0, 1} => 2, {0, 2} => 3, {1, 0} => 4, {1, 1} => 5 {1, 2} => 6 L'ordre dans lequel ces éléments sont ensuite stockés en mémoire est définies par l'implémentation. Les constantes Tensor présentent les contraintes suivantes:

• (C1) has_syntax(tensor_literal, element_type(tensor_type)), où: <ph type="x-smartling-placeholder">
  </ph>
  • has_syntax(element_literal: Syntax, element_type: Type) = is_wellformed(element_literal, type).
  • has_syntax(tensor_literal: List, element_type: Type) = has_syntax(tensor_literal..., element_type).
• (C2) has_shape(tensor_literal, shape(tensor_type)), où: <ph type="x-smartling-placeholder">
  </ph>
  • has_shape(element_literal: Syntax, []) = true.
  • has_shape(tensor_literal: List, shape: List) = size(tensor_literal) = shape[0] and has_shape(tensor_literal..., shape[1:]).
  • sinon false.

QuantizedTensorConstant ::= QuantizedTensorLiteral ':' QuantizedTensorType
QuantizedTensorLiteral  ::= 'dense' '<' (DenseLiteral | ElementLiteral) '>'

Les constantes de Tensors quantifiées représentent les valeurs de Tensor quantifiées à l'aide des mêmes en tant que constantes de Tensor, les éléments étant spécifiés comme des constantes de leur type de stockage. Les constantes de Tensor quantifiées présentent les contraintes suivantes:

• (C1) has_syntax(quantized_tensor_literal, storage_type(quantized_tensor_type)).
• (C2) has_shape(quantized_tensor_literal, shape(quantized_tensor_type)).

StringConstant  ::= StringLiteral
StringLiteral   ::= '"' {stringCharacter | escapeSequence} '"'
stringCharacter ::= all ASCII characters except '\00', '\01', ... '\1f' and '"'
escapeSequence  ::= '\' ('"' | '\' | 'n' | 't' | (hexadecimalDigit hexadecimalDigit))

Les littéraux de chaîne sont composés d'octets spécifiés à l'aide de caractères ASCII et et d'échappement. Comme ils sont indépendants de l'encodage, leur interprétation octets est défini par l'implémentation. Les littéraux de chaîne sont de type string.

Opérations

abs

Sémantique

Effectue une opération des abscisses par élément sur le Tensor operand et génère une result. Tensor. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les entiers signés: module entier.
• Pour les nombres à virgule flottante: abs d'IEEE-754.
• Pour les nombres complexes: module complexe.
• Pour les types quantifiés: dequantize_op_quantize(abs, operand, type(result)).

Entrées

Sorties

Contraintes

• (C1) shape(result) = shape(operand).
• (C2) baseline_element_type(result) est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • complex_element_type(element_type(operand)) si is_complex(operand).
  • baseline_element_type(operand) dans les autres cas.

Exemples

// %operand: [-2, 0, 2]
%result = "stablehlo.abs"(%operand) : (tensor<3xi32>) -> tensor<3xi32>
// %result: [2, 0, 2]

Autres exemples

add

Sémantique

Effectue l'addition par élément de deux Tensors lhs et rhs, et génère une Tensor result. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les valeurs booléennes: opérateur logique OU.
• Pour les entiers: addition d'entiers.
• Pour les nombres à virgule flottante: addition d'IEEE-754.
• Pour les nombres complexes: addition complexe.
• Pour les types quantifiés: dequantize_op_quantize(add, lhs, rhs, type(result)).

Entrées

Sorties

Contraintes

• Si l'opération utilise des Tensors non quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C1) type(lhs) = type(rhs) = type(result).
• Si l'opération utilise des Tensors quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C2) is_quantized(lhs) and is_quantized(rhs) and is_quantized(result).
  • (C3) storage_type(lhs) = storage_type(rhs) = storage_type(result).
  • (C4) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C5) (is_per_axis_quantized(lhs) or is_per_axis_quantized(rhs)) = is_per_axis_quantized(result).
  • (C6) Si is_per_axis_quantized(lhs), alors quantization_dimension(lhs) = quantization_dimension(result).
  • (C7) Si is_per_axis_quantized(rhs), alors quantization_dimension(rhs) = quantization_dimension(result).

Exemples

// %lhs: [[1, 2], [3, 4]]
// %rhs: [[5, 6], [7, 8]]
%result = "stablehlo.add"(%lhs, %rhs) : (tensor<2x2xi32>, tensor<2x2xi32>) -> tensor<2x2xi32>
// %result: [[6, 8], [10, 12]]

Autres exemples

after_all

Sémantique

Assurez-vous que les opérations à l'origine de l'inputs sont exécutées avant toute opérations qui dépendent de result. L'exécution de cette opération n'a aucun effet, il n'existe que pour établir les dépendances de données de result à inputs.

Entrées

Sorties

Exemples

// %input0: !stablehlo.token
// %input1: !stablehlo.token
%result = "stablehlo.after_all"(%input0, %input1) : (!stablehlo.token, !stablehlo.token) -> !stablehlo.token

Autres exemples

all_gather

Sémantique

Dans chaque groupe de processus de la grille de processus StableHLO, concatène les valeurs des Tensors operands de chaque processus le long de all_gather_dim et produit results.

L'opération divise la grille de processus StableHLO en process_groups, qui est définis comme suit:

• cross_replica(replica_groups) si channel_id <= 0 and use_global_device_ids = false.
• cross_replica_and_partition(replica_groups) si channel_id > 0 and use_global_device_ids = false.
• flattened_ids(replica_groups) si channel_id > 0 and use_global_device_ids = true.

Ensuite, dans chaque process_group:

• operands...@receiver = [operand@sender for sender in process_group] pour tous receiver dans process_group.
• results...@process = concatenate(operands...@process, all_gather_dim) pour tous process dans process_group.

Entrées

Sorties

Contraintes

• (C1) 0 <= all_gather_dim < rank(operands...).
• (C2) is_unique(replica_groups).
• (C3) size(replica_groups) est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas si cross_replica est utilisé.
  • num_replicas si cross_replica_and_partition est utilisé.
  • num_processes si flattened_ids est utilisé.
• (C4) 0 <= replica_groups < size(replica_groups).
• (C5) Si use_global_device_ids = true, alors channel_id > 0.
• (C6) type(results...) = type(operands...), sauf: <ph type="x-smartling-placeholder">
  </ph>
  • dim(results..., all_gather_dim) = dim(operands..., all_gather_dim) * dim(process_groups, 1).

Exemples

// num_replicas: 2
// num_partitions: 1
// %operand0@(0, 0): [[1, 2], [3, 4]]
// %operand0@(1, 0): [[5, 6], [7, 8]]
// %operand1@(0, 0): [[11, 12], [13, 14]]
// %operand1@(1, 0): [[15, 16], [17, 18]]
%result:2 = "stablehlo.all_gather"(%operand0, %operand1) {
  all_gather_dim = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  // channel_id = 0
  channel_handle = #stablehlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<2x2xi64>, tensor<2x2xi64>) -> (tensor<2x4xi64>, tensor<2x4xi64>)
// %result0@(0, 0): [[1, 2, 5, 6], [3, 4, 7, 8]]
// %result0@(1, 0): [[1, 2, 5, 6], [3, 4, 7, 8]]
// %result1@(0, 0): [[11, 12, 15, 16], [13, 14, 17, 18]]
// %result1@(1, 0): [[11, 12, 15, 16], [13, 14, 17, 18]]

Autres exemples

all_reduce

Sémantique

Dans chaque groupe de processus de la grille de processus StableHLO, applique une réduction fonction computation sur les valeurs des Tensors operands de chaque processus et produit des Tensors results.

L'opération divise la grille de processus StableHLO en process_groups, qui est définis comme suit:

• cross_replica(replica_groups) si channel_id <= 0 and use_global_device_ids = false.
• cross_replica_and_partition(replica_groups) si channel_id > 0 and use_global_device_ids = false.
• flattened_ids(replica_groups) si channel_id > 0 and use_global_device_ids = true.

Ensuite, dans chaque process_group:

• results...@process[result_index] = exec(schedule) pour une arborescence binaire schedule où: <ph type="x-smartling-placeholder">
  </ph>
  • exec(node) = computation(exec(node.left), exec(node.right)).
  • exec(leaf) = leaf.value.
• schedule est une arborescence binaire définie par l'implémentation, dont l'ordre le balayage est to_destination_type(operands...@process_group...[result_index], type(func_inputs(computation)[0])).

Entrées

Sorties

Contraintes

• (C1) is_unique(replica_groups).
• (C2) size(replica_groups) est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas si cross_replica est utilisé.
  • num_replicas si cross_replica_and_partition est utilisé.
  • num_processes si flattened_ids est utilisé.
• (C3) 0 <= replica_groups < size(replica_groups).
• (C4) Si use_global_device_ids = true, alors channel_id > 0.
• (C5) computation est de type (tensor<E>, tensor<E>) -> (tensor<E>), où is_promotable(element_type(operand), E)
• (C6) shape(results...) = shape(operands...).
• (C7) element_type(results...) = E.

Exemples

// num_replicas: 2
// num_partitions: 1
// %operand0@(0, 0): [1, 2, 3, 4]
// %operand0@(1, 0): [5, 6, 7, 8]
// %operand1@(0, 0): [9, 10, 11, 12]
// %operand1@(1, 0): [13, 14, 15, 16]
%result:2 = "stablehlo.all_reduce"(%operand0, %operand0) ({
  ^bb0(%arg0: tensor<i64>, %arg1: tensor<i64>):
    %0 = "stablehlo.add"(%arg0, %arg1) : (tensor<i64>, tensor<i64>) -> tensor<i64>
    "stablehlo.return"(%0) : (tensor<i64>) -> ()
}) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  // channel_id = 0
  channel_handle = #stablehlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<4xi64>, tensor<4xi64>) -> (tensor<4xi64>, tensor<4xi64>)
// %result0@(0, 0): [6, 8, 10, 12]
// %result0@(1, 0): [6, 8, 10, 12]
// %result1@(0, 0): [22, 24, 26, 28]
// %result1@(1, 0): [22, 24, 26, 28]

Autres exemples

all_to_all

Sémantique

Dans chaque groupe de processus de la grille de processus StableHLO, divise les valeurs de les Tensors operands le long de split_dimension en plusieurs parties, dispersent la division entre les processus, concatène les parties dispersées concat_dimension et génère des Tensors results. L'opération divise la grille de processus StableHLO en process_groups, qui est définis comme suit:

• cross_replica(replica_groups) si channel_id <= 0.
• cross_partition(replica_groups) si channel_id > 0.

Ensuite, dans chaque process_group:

• split_parts...@sender = split(operands...@sender, split_count, split_dimension) pour les sender de process_group.
• scattered_parts...@receiver = [split_parts...@sender[receiver_index] for sender in process_group] où receiver_index = process_group.index(receiver)
• results...@process = concatenate(scattered_parts...@process, concat_dimension).

Entrées

Sorties

Contraintes

• (C1) 0 <= split_dimension < rank(operands...).
• (C2) dim(operands..., split_dimension) % split_count = 0.
• (C3) 0 <= concat_dimension < rank(operands...).
• (C4) 0 < split_count.
• (C5) is_unique(replica_groups).
• (C6) size(replica_groups) est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas si cross_replica est utilisé.
  • num_partitions si cross_partition est utilisé.
• (C7) 0 <= replica_groups < size(replica_groups).
• (C8) dim(replica_groups, 1) = split_count.
• (C9) type(results...) = type(operands...) sauf, si split_dimension != concat_dimension: <ph type="x-smartling-placeholder">
  </ph>
  • dim(results..., split_dimension) = dim(operands..., split_dimension) / split_count.
  • dim(results..., concat_dimension) = dim(operands..., concat_dimension) * split_count.

Exemples

// num_replicas: 2
// num_partitions: 1
// %operand1@(0, 0): [[1, 2, 3, 4],
//                    [5, 6, 7, 8]]
// %operand1@(1, 0): [[9, 10, 11, 12],
//                    [13, 14, 15, 16]]
// %operand2@(0, 0): [[17, 18, 19, 20],
//                    [21, 22, 23, 24]]
// %operand2@(1, 0): [[25, 26, 27, 28],
//                    [29, 30, 31, 32]]
%result:2 = "stablehlo.all_to_all"(%operand1, %operand2) {
  split_dimension = 1 : i64,
  concat_dimension = 0 : i64,
  split_count = 2 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
} : (tensor<2x4xi64>, tensor<2x4xi64>) -> (tensor<4x2xi64>, tensor<4x2xi64>)
// %result#0@(0, 0): [[1, 2], [5, 6], [9, 10], [13, 14]]
// %result#0@(1, 0): [[3, 4], [7, 8], [11, 12], [15, 16]]
// %result#1@(0, 0): [[17, 18], [21, 22], [25, 26], [29, 30]]
// %result#1@(1, 0): [[19, 20], [23, 24], [27, 28], [31, 32]]

Autres exemples

et

Sémantique

Effectue l'opérateur AND par élément de deux Tensors lhs et rhs et génère une result Tensor. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les valeurs booléennes: l'opérateur logique AND.
• Pour les entiers: AND (ET) bit à bit.

Entrées

Sorties

Contraintes

• (C1) type(lhs) = type(rhs) = type(result).

Exemples

// %lhs: [[1, 2], [3, 4]]
// %rhs: [[5, 6], [7, 8]]
%result = "stablehlo.and"(%lhs, %rhs) : (tensor<2x2xi32>, tensor<2x2xi32>) -> tensor<2x2xi32>
// %result: [[1, 2], [3, 0]]

Autres exemples

atan2

Sémantique

Effectue des opérations atan2 au niveau des éléments sur les Tensors lhs et rhs, et génère une Tensor result. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les nombres à virgule flottante: atan2 d'IEEE-754.
• Pour les nombres complexes: complexe atan2.
• Pour les types quantifiés: dequantize_op_quantize(atan2, lhs, rhs, type(result)).

Entrées

Sorties

Contraintes

• (C1) baseline_type(lhs) = baseline_type(rhs) = baseline_type(result).

Exemples

// %lhs: [0.0, 1.0, -1.0]
// %rhs: [0.0, 0.0, 0.0]
%result = "stablehlo.atan2"(%lhs, %rhs) : (tensor<3xf64>, tensor<3xf64>) -> tensor<3xf64>
// %result: [0.0, 1.57079637, -1.57079637] // [0.0, pi/2, -pi/2]

Autres exemples

batch_norm_grad

Sémantique

Calcule les gradients de plusieurs entrées de la rétropropagation de batch_norm_training. à partir de grad_output, et génère grad_operand, grad_scale et grad_offset Tensors. Plus formellement, cette opération peut être exprimée sous la forme d'une décomposition Opérations StableHLO existantes à l'aide de la syntaxe Python, comme suit:

def compute_sum(operand, feature_index):
  (sum,) = reduce(
      inputs=[operand],
      init_values=[constant(0, element_type(operand))],
      dimensions=[i for i in range(rank(operand)) if i != feature_index],
      body=lambda x, y: add(x, y))
  return sum

def compute_mean(operand, feature_index):
  sum = compute_sum(operand, feature_index)
  divisor = constant(size(operand) / dim(operand, feature_index),
                     element_type(operand))
  divisor_bcast = broadcast_in_dim(divisor, [], type(sum))
  return divide(sum, divisor_bcast)

def batch_norm_grad(operand, scale, mean, variance, grad_output, epsilon, feature_index):
  # Broadcast inputs to type(operand)
  scale_bcast = broadcast_in_dim(scale, [feature_index], type(operand))
  mean_bcast = broadcast_in_dim(mean, [feature_index], type(operand))
  variance_bcast = broadcast_in_dim(variance, [feature_index], type(operand))
  epsilon_bcast = broadcast_in_dim(constant(epsilon, element_type(operand)), [],
                                   type(operand))

  # Perform normalization using the provided `mean` and `variance`
  # Intermediate values will be useful for computing gradients
  centered_operand = subtract(operand, mean_bcast)
  stddev = sqrt(add(variance_bcast, epsilon_bcast))
  normalized_operand = divide(centered_operand, stddev)

  # Use the implementation from batchnorm_expander.cc in XLA
  # Temporary variables have exactly the same names as in the C++ code
  elements_per_feature = broadcast_in_dim(
      constant(divide(size(operand), dim(operand, feature_index)),
               element_type(grad_output)),
      [], type(operand))
  i1 = multiply(grad_output, elements_per_feature)
  i2 = broadcast_in_dim(
      compute_sum(grad_output, feature_index), [feature_index], type(operand))
  i3 = broadcast_in_dim(
      compute_sum(multiply(grad_output, centered_operand), feature_index),
      [feature_index], type(operand))
  i4 = multiply(i3, centered_operand)
  i5 = divide(i4, add(variance_bcast, epsilon_bcast))
  i6 = subtract(subtract(i1, i2), i5)

  grad_operand =
      multiply(divide(divide(scale_bcast, stddev), elements_per_feature), i6)
  grad_scale =
      compute_sum(multiply(grad_output, normalized_operand), feature_index)
  grad_offset = compute_sum(grad_output, feature_index)

  return grad_operand, grad_scale, grad_offset

Pour les types quantifiés, exécute dequantize_batch_norm_grad_or_training_quantize(lambda operand, scale, mean, variance, grad_output: batch_norm_grad(operand, scale, mean, variance, grad_output, epsilon, feature_index), operand, scale, mean, variance, grad_output, type(grad_operand), type(grad_scale), type(feature_index))

Entrées

Sorties

Contraintes

• (C1) 0 <= feature_index < rank(operand).
• (C2) operand, scale, mean, variance, grad_output, grad_operand grad_scale et grad_offset ont le même baseline_element_type.
• (C3) operand, grad_output et grad_operand ont la même forme.
• (C4) scale, mean, variance, grad_scale et grad_offset ont même forme.
• (C5) size(scale) = dim(operand, feature_index).

Exemples

// %operand: [
//            [[1.0, 2.0], [3.0, 4.0]],
//            [[3.0, 4.0], [1.0, 2.0]]
//           ]
// %scale: [1.0, 1.0]
// %mean: [2.0, 3.0]
// %variance: [1.0, 1.0]
// %grad_output: [
//                [[0.1, 0.1], [0.1, 0.1]],
//                [[0.1, 0.1], [0.1, 0.1]]
//               ]
%grad_operand, %grad_scale, %grad_offset =
"stablehlo.batch_norm_grad"(%operand, %scale, %mean, %variance, %grad_output) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf64>, tensor<2xf64>, tensor<2xf64>, tensor<2xf64>,
     tensor<2x2x2xf64>) -> (tensor<2x2x2xf64>, tensor<2xf64>, tensor<2xf64>)
// %grad_operand: [
//                 [[0.0, 0.0], [0.0, 0.0]],
//                 [[0.0, 0.0], [0.0, 0.0]]
//                ]
// %grad_scale:  [0.0, 0.0]
// %grad_offset: [0.4, 0.4]

batch_norm_inference

Sémantique

Normalise le Tensor operand pour toutes les dimensions, à l'exception du feature_index et génère un Tensor result. Plus formellement, cela peut être exprimée sous la forme d'une décomposition en opérations StableHLO existantes à l'aide de la syntaxe Python suivante:

def batch_norm_inference(operand, scale, offset, mean, variance, epsilon, feature_index):
  # Broadcast inputs to shape(operand)
  scale_bcast = broadcast_in_dim(scale, [feature_index], type(operand))
  offset_bcast = broadcast_in_dim(offset, [feature_index], type(operand))
  mean_bcast = broadcast_in_dim(mean, [feature_index], type(operand))
  variance_bcast = broadcast_in_dim(variance, [feature_index], type(operand))
  epsilon_bcast = broadcast_in_dim(constant(epsilon, element_type(operand)), [],
                                   type(operand))

  # Perform normalization using the provided `mean` and `variance` instead of
  # computing them like `batch_norm_training` does.
  centered_operand = subtract(operand, mean_bcast)
  stddev = sqrt(add(variance_bcast, epsilon_bcast))
  normalized_operand = divide(centered_operand, stddev)
  return add(multiply(scale_bcast, normalized_operand), offset_bcast)

Pour les types quantifiés, exécute dequantize_op_quantize(lambda operand, scale, offset, mean, variance: batch_norm_inference(operand, scale, offset, mean, variance, epsilon, feature_index), operand, scale, offset, mean, variance, type(result))

Entrées

Sorties

Contraintes

• (C1) 0 <= feature_index < rank(operand).
• (C2) operand, scale, offset, mean, variance et result ont même baseline_element_type.
• (C3) size(scale) = dim(operand, feature_index).
• (C4) size(offset) = dim(operand, feature_index).
• (C5) size(mean) = dim(operand, feature_index).
• (C6) size(variance) = dim(operand, feature_index).
• (C7) baseline_type(operand) = baseline_type(result).

Exemples

// %operand: [
//            [[1.0, 2.0], [3.0, 4.0]],
//            [[3.0, 4.0], [1.0, 2.0]]
//           ]
// %scale: [1.0, 1.0]
// %offset: [1.0, 1.0]
// %mean: [2.0, 3.0]
// %variance: [1.0, 1.0]
%result = "stablehlo.batch_norm_inference"(%operand, %scale, %offset, %mean, %variance) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf64>, tensor<2xf64>, tensor<2xf64>, tensor<2xf64>, tensor<2xf64>) -> tensor<2x2x2xf64>
// %result: [
//           [[0.0, 0.0], [2.0, 2.0]],
//           [[2.0, 2.0], [0.0, 0.0]]
//          ]

batch_norm_training

Sémantique

Calcule la moyenne et la variance pour toutes les dimensions, à l'exception de feature_index et normalise le Tensor operand générant output, batch_mean et batch_var. Plus formellement, cette opération peut être exprimée sous la forme décomposition en opérations StableHLO existantes à l'aide de la syntaxe Python en tant que ce qui suit:

def compute_mean(operand, feature_index):
  (sum,) = reduce(
      inputs=[operand],
      init_values=[constant(0, element_type(operand))],
      dimensions=[i for i in range(rank(operand)) if i != feature_index],
      body=lambda x, y: add(x, y))
  divisor = constant(size(operand) / dim(operand, feature_index),
                     element_type(operand))
  divisor_bcast = broadcast_in_dim(divisor, [], type(sum))
  return divide(sum, divisor_bcast)

def compute_variance(operand, feature_index):
  mean = compute_mean(operand, feature_index)
  mean_bcast = broadcast_in_dim(mean, [feature_index], type(operand))
  centered_operand = subtract(operand, mean_bcast)
  return compute_mean(mul(centered_operand, centered_operand), feature_index)

def batch_norm_training(operand, scale, offset, epsilon, feature_index):
  mean = compute_mean(operand, feature_index)
  variance = compute_variance(operand, feature_index)
  return batch_norm_inference(operand, scale, offset, mean, variance, epsilon,
                              feature_index),
         mean, variance

Pour les types quantifiés, exécute dequantize_batch_norm_grad_or_training_quantize(lambda operand, scale, offset: batch_norm_training(operand, scale, offset, epsilon, feature_index), operand, scale, offset, type(output), type(batch_mean), type(batch_var))

Entrées

Sorties

Contraintes

• (C1) 0 <= feature_index < rank(operand).
• (C2) operand, scale, offset, batch_mean, batch_var et output ont le même baseline_element_type.
• (C3) size(scale) = dim(operand, feature_index).
• (C4) size(offset) = dim(operand, feature_index).
• (C5) size(batch_mean) = dim(operand, feature_index).
• (C6) size(batch_var) = dim(operand, feature_index).
• (C7) baseline_type(output) = baseline_type(operand).

Exemples

// %operand: [
//            [[1.0, 2.0], [3.0, 4.0]],
//            [[3.0, 4.0], [1.0, 2.0]]
//           ]
// %scale: [1.0, 1.0]
// %offset: [1.0, 1.0]
%output, %batch_mean, %batch_var = "stablehlo.batch_norm_training"(%operand, %scale, %offset) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf64>, tensor<2xf64>, tensor<2xf64>) ->
    (tensor<2x2x2xf64>, tensor<2xf64>, tensor<2xf64>)
// %output: [
//           [[0.0, 0.0], [2.0, 2.0]],
//           [[2.0, 2.0], [0.0, 0.0]]
//          ]
// %batch_mean: [2.0, 3.0]
// %batch_var: [1.0, 1.0]

bitcast_convert

Sémantique

Effectue une opération de diffusion de bits sur le Tensor operand et génère un Tensor result. où les bits de l'ensemble du Tensor operand sont réinterprétés à l'aide du Type de Tensor result.

Plus formellement, étant donné que E = element_type(operand), E' = element_type(result), et R = rank(operand):

• Si la valeur est num_bits(E') < num_bits(E), bits(result[i0, ..., iR-1, :]) = bits(operand[i0, ..., iR-1])
• Si la valeur est num_bits(E') > num_bits(E), bits(result[i0, ..., iR-2]) = bits(operand[i0, ..., iR-2, :])
• Si la valeur est num_bits(E') = num_bits(E), bits(result[i0, ..., iR-1]) = bits(operand[i0, ..., iR-1])

bits renvoie la représentation en mémoire d'une valeur donnée et de son comportement. est définie par l'implémentation, car la représentation exacte des Tensors est définie par l'implémentation, et la représentation exacte des types d'éléments définies par l'implémentation.

Entrées

Sorties

Contraintes

• (C1) Avec E = is_quantized(operand) ? storage_type(operand) : element_type(operand), E' = is_quantized(result) ? storage_type(result) : element_type(result) et R = rank(operand): <ph type="x-smartling-placeholder">
  </ph>
  • Si num_bits(E') = num_bits(E), shape(result) = shape(operand).
  • Si la valeur est num_bits(E') < num_bits(E):
  • rank(result) = R + 1.
  • dim(result, i) = dim(operand, i) pour tous les 0 <= i < R.
  • dim(result, R) * num_bits(E') = num_bits(E).
  • Si la valeur est num_bits(E') > num_bits(E):
  • rank(result) = R - 1.
  • dim(result, i) = dim(operand, i) pour tous les 0 <= i < R.
  • dim(operand, R - 1) * num_bits(E) = num_bits(E').
• (C2) Si is_complex(operand) or is_complex(result), alors is_complex(operand) and is_complex(result)

Exemples

// %operand: 0x0123456789ABCDEF
%result = "stablehlo.bitcast_convert"(%operand) : (tensor<f64>) -> tensor<4xf16>
// %result: [0xCDEF, 0x89AB, 0x4567, 0x0123] // little-endian representation

Autres exemples

broadcast_in_dim

Sémantique

Développe les dimensions et/ou le rang d'un Tensor d'entrée en dupliquant les données dans le Tensor operand et produit un Tensor result. Plus formellement, result[result_index] = operand[operand_index] où pour les d de axes(operand):

• operand_index[d] = 0 si dim(operand, d) = 1.
• operand_index[d] = result_index[broadcast_dimensions[d]] dans les autres cas.

Entrées

Sorties

Contraintes

• (C1) La valeur element_type(result) est donnée par: <ph type="x-smartling-placeholder">
  </ph>
  • element_type(operand), si !is_per_axis_quantized(operand).
  • element_type(operand), à l'exception de quantization_dimension(operand), scales(operand) et zero_points(operand) peuvent être différents de quantization_dimension(result), scales(result) et zero_points(result) resp., sinon.
• (C2) size(broadcast_dimensions) = rank(operand).
• (C3) 0 <= broadcast_dimensions < rank(result).
• (C4) is_unique(broadcast_dimensions).
• (C5) Pour tous les d de axes(operand): <ph type="x-smartling-placeholder">
  </ph>
  • dim(operand, d) = 1 ou
  • dim(operand, d) = dim(result, broadcast_dimensions[d]).
• (C6) Si is_per_axis_quantized(result): <ph type="x-smartling-placeholder">
  </ph>
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)].
  • Si la valeur est dim(operand, quantization_dimension(operand)) = 1, alors scales(result)[i] = scales(operand)[0] and zero_points(result)[i] = zero_points(operand)[0] for i in range(dim(result, quantization_dimension(result)))

Exemples

// %operand: [
//            [1, 2, 3]
//           ]
%result = "stablehlo.broadcast_in_dim"(%operand) {
  broadcast_dimensions = array<i64: 2, 1>
} : (tensor<1x3xi32>) -> tensor<2x3x2xi32>
// %result: [
//            [
//             [1, 1],
//             [2, 2],
//             [3, 3]
//            ],
//            [
//             [1, 1],
//             [2, 2],
//             [3, 3]
//            ]
//          ]

Autres exemples

coque

Sémantique

Génère le résultat de l'exécution d'une seule fonction de branches en fonction de la valeur de index. Plus formellement, result = selected_branch() où:

• selected_branch = branches[index] si 0 <= index < size(branches).
• selected_branch = branches[-1] dans les autres cas.

Entrées

Sorties

Contraintes

• (C1) 0 < size(branches).
• (C2) input_types(branches...) = [].
• (C3) same(output_types(branches...)).
• (C4) type(results...) = output_types(branches[0]).

Exemples

// %index: -1
// %result_branch0: [0, 0]
// %result_branch1: [1, 1]
%result0, %result1 = "stablehlo.case"(%index) ({
  "stablehlo.return"(%result_branch0, %result_branch0) : (tensor<2xi64>, tensor<2xi64>) -> ()
}, {
  "stablehlo.return"(%result_branch1, %result_branch1) : (tensor<2xi64>, tensor<2xi64>) -> ()
}) : (tensor<i32>) -> (tensor<2xi64>, tensor<2xi64>)
// %result0: [1, 1]
// %result1: [1, 1]

Autres exemples

cbrt

Sémantique

Effectue une opération racine cubique par élément sur le Tensor operand et génère une Tensor result. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les nombres à virgule flottante: rootn(x, 3) d'IEEE-754.
• Pour les nombres complexes: racine cubique complexe.
• Pour les types quantifiés: dequantize_op_quantize(cbrt, operand, type(result))

Entrées

Sorties

Contraintes

• (C1) baseline_type(operand) = baseline_type(result).

Exemples

// %operand: [0.0, 1.0, 8.0, 27.0]
%result = "stablehlo.cbrt"(%operand) : (tensor<4xf64>) -> tensor<4xf64>
// %result: [0.0, 1.0, 2.0, 3.0]

Autres exemples

ceil

Sémantique

Effectue un ceil par élément du Tensor operand et produit un Tensor result. Elle met en œuvre l'opération roundToIntegralTowardPositive de la norme IEEE-754. spécifique. Pour les types quantifiés, exécute dequantize_op_quantize(ceil, operand, type(result))

Entrées

Sorties

Contraintes

• (C1) baseline_type(operand) = baseline_type(result).

Exemples

// %operand: [-0.8166, -0.2530, 0.2530, 0.8166, 2.0]
%result = "stablehlo.ceil"(%operand) : (tensor<5xf32>) -> tensor<5xf32>
// %result: [-0.0, -0.0, 1.0, 1.0, 2.0]

Autres exemples

Cholesky

Sémantique

Calcule la décomposition par Cholesky d'un lot de matrices.

Plus formellement, pour tous les i de index_space(result), result[i0, ..., iR-3, :, :] est une décomposition cholesky de a[i0, ..., iR-3, :, :], qui se présente sous la forme d'un triangle inférieur (si lower est true) ou une matrice triangulaire supérieure (si lower est false). Les valeurs de sortie dans le triangle opposé, c'est-à-dire le triangle supérieur strict le triangle inférieur strict, sont définies par l'implémentation.

S'il existe des valeurs i où la matrice d'entrée n'est pas de type hermitien , le comportement n'est pas défini.

Pour les types quantifiés, exécute dequantize_op_quantize(lambda operand: cholesky(operand, lower), a, type(result))

Entrées

Sorties

Contraintes

• (C1) baseline_type(a) = baseline_type(result).
• (C2) 2 <= rank(a).
• (C3) dim(a, -2) = dim(a, -1).

Exemples

// %a: [
//      [1.0, 2.0, 3.0],
//      [2.0, 20.0, 26.0],
//      [3.0, 26.0, 70.0]
//     ]
%result = "stablehlo.cholesky"(%a) {
  lower = true
} : (tensor<3x3xf32>) -> tensor<3x3xf64>
// %result: [
//           [1.0, 0.0, 0.0],
//           [2.0, 4.0, 0.0],
//           [3.0, 5.0, 6.0]
//          ]

limiter

Sémantique

Bloque chaque élément du Tensor operand entre une valeur minimale et une valeur maximale et produit un Tensor result. Plus formellement, result[result_index] = minimum(maximum(operand[result_index], min_element), max_element), où min_element = rank(min) = 0 ? min[] : min[result_index], max_element = rank(max) = 0 ? max[] : max[result_index] Pour les types quantifiés, exécute dequantize_op_quantize(clamp, min, operand, max, type(result)).

Ordonner des nombres complexes implique une sémantique surprenante, C'est pourquoi nous prévoyons de ne plus prendre en charge les nombres complexes à l'avenir. pour cette opération (#560).

Entrées

Sorties

Contraintes

• (C1) rank(min) = 0 or shape(min) = shape(operand).
• (C2) rank(max) = 0 or shape(max) = shape(operand).
• (C3) baseline_element_type(min) = baseline_element_type(operand) = baseline_element_type(max).
• (C4) baseline_type(operand) = baseline_type(result).

Exemples

// %min: [5, 10, 15]
// %operand: [3, 13, 23]
// %max: [10, 15, 20]
%result = "stablehlo.clamp"(%min, %operand, %max) : (tensor<3xi32>, tensor<3xi32>, tensor<3xi32>) -> tensor<3xi32>
// %result: [5, 13, 20]

Autres exemples

collective_broadcast

Sémantique

Dans chaque groupe de processus de la grille de processus StableHLO, envoyez la valeur du Tensor operand du processus source aux processus cibles et produit une Tensor result.

L'opération divise la grille de processus StableHLO en process_groups, qui est définis comme suit:

• cross_replica(replica_groups) si channel_id <= 0.
• cross_partition(replica_groups) si channel_id > 0.

Ensuite, result@process est donné par:

• operand@process_groups[i, 0] s'il existe un i de sorte que le processus soit dans le pays suivant : process_groups[i].
• broadcast_in_dim(constant(is_quantized(result) ? quantize(0, element_type(result)) : 0, element_type(result)), [], type(result)) sinon.

Entrées

Sorties

Contraintes

• (C1) is_unique(replica_groups).
• (C2) 0 <= replica_groups < N, où N est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas si cross_replica est utilisé.
  • num_partitions si cross_partition est utilisé.
• (C3) type(result) = type(operand).

Exemples

// num_replicas: 4
// num_partitions: 1
// %operand@(0, 0): [[1, 2]]
// %operand@(1, 0): [[3, 4]]
// %operand@(2, 0): [[5, 6]]
// %operand@(3, 0): [[7, 8]]
%result = "stablehlo.collective_broadcast"(%operand) {
  replica_groups = dense<[[2, 1]]> : tensor<1x2xi64>,
  channel_handle = #stablehlo.channel_handle<handle = 0, type = 0>
} : (tensor1x2xi64>) -> tensor<1x2xi64>
// %result@(0, 0): [[0, 0]]
// %result@(1, 0): [[5, 6]]
// %result@(2, 0): [[5, 6]]
// %result@(3, 0): [[0, 0]]

collective_permute

Sémantique

Dans chaque groupe de processus de la grille de processus StableHLO, envoie la valeur du Tensor operand du processus source au processus cible et produit une Tensor result.

L'opération divise la grille de processus StableHLO en process_groups, qui est définis comme suit:

• cross_replica(source_target_pairs) si channel_id <= 0.
• cross_partition(source_target_pairs) si channel_id > 0.

Ensuite, result@process est donné par:

• operand@process_groups[i, 0], s'il existe un i tel que process_groups[i, 1] = process
• broadcast_in_dim(constant(is_quantized(result) ? quantize(0, element_type(result)) : 0, element_type(result)), [], type(result)) sinon.

Entrées

Sorties

Contraintes

• (C1) dim(source_target_pairs, 1) = 2.
• (C2) is_unique(source_target_pairs[:, 0]).
• (C3) is_unique(source_target_pairs[:, 1]).
• (C4) 0 <= source_target_pairs < N, où N est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas si cross_replica est utilisé.
  • num_partitions si cross_partition est utilisé.
• (C5) type(result) = type(operand).

Exemples

// num_replicas: 3
// num_partitions: 1
// %operand@(0, 0): [[1, 2], [3, 4]]
// %operand@(1, 0): [[5, 6], [7, 8]]
// %operand@(2, 0): [[9, 10], [11, 12]]
%result = "stablehlo.collective_permute"(%operand) {
  source_target_pairs = dense<[[0, 1], [1, 2]]> : tensor<2x2xi64>,
  channel_handle = #stablehlo.channel_handle<handle = 0, type = 0>
} : (tensor<2x2xi64>) -> tensor<2x2xi64>
//
// %result@(0, 0): [[0, 0], [0, 0]]
// %result@(1, 0): [[1, 2], [3, 4]]
// %result@(2, 0): [[5, 6], [7, 8]]

Autres exemples

compare

Sémantique

Effectue une comparaison par élément des Tensors lhs et rhs en fonction comparison_direction et compare_type, et génère un Tensor result.

Les valeurs de comparison_direction et compare_type ont les valeurs suivantes : sémantique:

Pour les éléments de type booléen et entier:

• EQ : lhs = rhs.
• NE : lhs != rhs.
• GE : lhs >= rhs.
• GT : lhs > rhs.
• LE : lhs <= rhs.
• LT : lhs < rhs.

Pour les types d'éléments à virgule flottante avec compare_type = FLOAT, l'opération implémente les opérations IEEE-754 suivantes:

• EQ : compareQuietEqual.
• NE : compareQuietNotEqual.
• GE : compareQuietGreaterEqual.
• GT : compareQuietGreater.
• LE : compareQuietLessEqual.
• LT : compareQuietLess.

Pour les types d'éléments à virgule flottante avec compare_type = TOTALORDER, l'opération utilise la combinaison des opérations totalOrder et compareQuietEqual de IEEE-754.

Pour les types d'éléments complexes, la comparaison lexicographique des paires (real, imag) est effectuées à l'aide des méthodes comparison_direction et compare_type fournies. Ordonner des nombres complexes implique une sémantique surprenante, C'est pourquoi nous prévoyons de ne plus prendre en charge les nombres complexes à l'avenir. lorsque comparison_direction est GE, GT, LE ou LT (560)

Pour les types quantifiés. exécute dequantize_compare(lhs, rhs, comparison_direction).

Entrées

Sorties

Contraintes

• (C1) baseline_element_type(lhs) = baseline_element_type(rhs).
• (C2) shape(lhs) = shape(rhs) = shape(result).
• (C3) compare_type est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • SIGNED si is_signed_integer(element_type(lhs)).
  • UNSIGNED si is_unsigned_integer(element_type(lhs)) or is_boolean(element_type(lhs)).
  • FLOAT ou TOTALORDER si is_float(element_type(lhs)).
  • FLOAT si is_complex(element_type(lhs)).

Exemples

// %lhs: [1.0, 3.0]
// %rhs: [1.1, 2.9]
%result = "stablehlo.compare"(%lhs, %rhs) {
  comparison_direction = #stablehlo<comparison_direction LT>,
  compare_type = #stablehlo<comparison_type FLOAT>
} : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>
// %result: [true, false]

Autres exemples

complexe

Sémantique

Effectue une conversion au niveau des éléments en une valeur complexe à partir d'une paire de valeurs réelles et des valeurs imaginaires, lhs et rhs, et produit un Tensor result.

Entrées

Sorties

Contraintes

• (C1) type(lhs) = type(rhs).
• (C2) shape(result) = shape(lhs).
• (C3) element_type(result) est de type complex<E>, où E = element_type(lhs)

Exemples

// %lhs: [1.0, 3.0]
// %rhs: [2.0, 4.0]
%result = "stablehlo.complex"(%lhs, %rhs) : (tensor<2xf64>, tensor<2xf64>) -> tensor<2xcomplex<f64>>
// %result: [(1.0, 2.0), (3.0, 4.0)]

Autres exemples

composite

Sémantique

Encapsule une opération composée d'autres opérations StableHLO. en prenant inputs et composite_attributes, et en produisant results. La de l'opération est implémentée par l'attribut decomposition. La L'opération composite peut être remplacée par sa décomposition sans modifier le programme la sémantique. Dans les cas où l'intégration de la décomposition ne fournit pas utilisez plutôt custom_call.

Le champ version (par défaut, 0) est utilisé pour indiquer quand un élément composite est changement de sémantique.

Entrées

Sorties

Contraintes

• (C1) is_namespaced_op_name(name)
• (C2) is_defined_in_parent_scope(decomposition)
• (C3) types(inputs...) == input_types(decomposition)
• (C4) types(results...) == output_types(decomposition)

Exemples

%results = "stablehlo.composite"(%input0, %input1) {
  name = "my_namespace.my_op",
  composite_attributes = {
    my_attribute = "my_value"
  },
  decomposition = @my_op,
  version = 1 : i32
} : (tensor<f32>, tensor<f32>) -> tensor<f32>

Autres exemples

concatenate

Sémantique

Concatène inputs avec la dimension dimension dans le même ordre que celui donné et génère un Tensor result. Plus formellement, result[i0, ..., id, ..., iR-1] = inputs[k][i0, ..., kd, ..., iR-1], où:

1. id = d0 + ... + dk-1 + kd.
2. d est égal à dimension, et d0... correspond à la de taille de dimension sur inputs.

Entrées

Sorties

Contraintes

• (C1) same(element_type(inputs...)).
• (C2) same(shape(inputs...)), sauf dim(inputs..., dimension).
• (C3) 0 < size(inputs).
• (C4) 0 <= dimension < rank(inputs[0]).
• (C5) element_type(result) = element_type(inputs[0]).
• (C6) shape(result) = shape(inputs[0]), sauf: <ph type="x-smartling-placeholder">
  </ph>
  • dim(result, dimension) = dim(inputs[0], dimension) + ....

Exemples

// %input0: [[1, 2], [3, 4], [5, 6]]
// %input1: [[7, 8]]
%result = "stablehlo.concatenate"(%input0, %input1) {
  dimension = 0 : i64
} : (tensor<3x2xi64>, tensor<1x2xi64>) -> tensor<4x2xi64>
// %result: [[1, 2], [3, 4], [5, 6], [7, 8]]

Autres exemples

constante

Sémantique

Génère un Tensor output à partir d'une constante value.

Entrées

Sorties

Contraintes

• (C1) type(value) = type(output).

Exemples

%output = "stablehlo.constant"() {
  value = dense<[[0.0, 1.0], [2.0, 3.0]]> : tensor<2x2xf32>
} : () -> tensor<2x2xf32>
// %output: [[0.0, 1.0], [2.0, 3.0]]

Autres exemples

d'effectuer une conversion

Sémantique

Effectue une conversion au niveau des éléments d'un type d'élément à un autre sur operand et produit un Tensor result.

Pour les conversions boolean-to-any-supported-type, la valeur false est convertie en zéro, et la valeur true en un. Pour any-supported-type-to-boolean, une valeur nulle est convertie en false et les valeurs non nulles sont converties en true. Découvrez ci-dessous comment cela fonctionnent pour des types complexes.

Pour les conversions impliquant entier en entier, entier en virgule flottante ou floating-point-to-floating-point, si la valeur source peut être exactement représentée dans le type "destination", la valeur du résultat correspond exactement représentation. Sinon, le comportement reste à déterminer (#180)

Pour les conversions impliquant une conversion floating-point-to-integer, la partie fractionnaire est la suivante : tronquées. Si la valeur tronquée ne peut pas être représentée dans le type de destination, comportement à déterminer (#180).

Les conversions complexes à complexes suivent le même comportement : les conversions floating-point-to-floating-point pour convertir des valeurs réelles et des parties imaginaires.

Pour les conversions de type complex-to-any-other-type et complex-to-any-other-type : la valeur imaginaire source est ignorée, ou la valeur imaginaire de destination est mis à zéro, respectivement. La conversion de la partie réelle suit les conversions à virgule flottante.

En principe, cette opération peut exprimer une déquantification (conversion à partir de Tensors quantifiés en Tensors standards), la quantification (conversion des Tensors en Tensors quantifiés) et requantification (conversion entre des Tensors quantifiés) mais pour le moment, nous avons des opérations dédiées. uniform_dequantize pour le premier cas d'utilisation et uniform_quantize pour le deuxième et troisième cas d'utilisation. À l'avenir, ces deux opérations pourront être fusionnées dans convert (#1576).

Entrées

Sorties

Contraintes

• (C1) shape(operand) = shape(result).

Exemples

// %operand: [-1, 0, 1]
%result = "stablehlo.convert"(%operand) : (tensor<3xi64>) -> tensor<3xcomplex<f64>>
// %result: [(-1.0, 0.0), (0.0, 0.0), (1.0, 0.0)]

Autres exemples

Convolution

Sémantique

Calcule les produits scalaires entre les fenêtres de lhs et des tranches de rhs, puis produit result Le schéma suivant montre comment les éléments de result sont calculés à partir de lhs et rhs à l'aide d'un exemple concret.

Plus formellement, envisagez de recadrer les entrées ci-dessous en termes de lhs. pour pouvoir exprimer des fenêtres de lhs:

• lhs_window_dimensions = lhs_shape(dim(lhs, input_batch_dimension), dim(rhs, kernel_spatial_dimensions), dim(lhs, input_feature_dimension)).
• lhs_window_strides = lhs_shape(1, window_strides, 1)
• lhs_padding = lhs_shape([0, 0], padding, [0, 0])
• lhs_base_dilations = lhs_shape(1, lhs_dilation, 1)
• lhs_window_dilations = lhs_shape(1, rhs_dilation, 1).

Ce recadrement utilise les fonctions d'assistance suivantes:

• lhs_shape(n, hw, c) = permute([n] + hw + [c], [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension]).
• result_shape(n1, hw, c1) = permute([n1] + hw + [c1], [output_batch_dimension] + output_spatial_dimensions + [output_feature_dimension]).
• permute([j0, j1, ..., jR-1], permutation) = [i0, i1, ..., iR-1], où j[d] = i[permutation[d]].

Si feature_group_count = 1 et batch_group_count = 1, alors pour toutes output_spatial_index à index_space(dim(result, output_spatial_dimensions...)), result[result_shape(:, output_spatial_index, :)] = dot_product, où:

• padding_value = constant(0, element_type(lhs)).
• padded_lhs = pad(lhs, padding_value, lhs_padding[:, 0], lhs_padding[:, 1], lhs_base_dilations - 1)
• lhs_window_start = lhs_shape(0, output_spatial_index, 0) * lhs_window_strides
• lhs_window = slice(padded_lhs, lhs_window_start, lhs_window_start + lhs_window_dimensions, lhs_window_dilations).
• reversed_lhs_window = reverse(lhs_window, [input_spatial_dimensions[dim] for dim in range(size(window_reversal)) if window_reversal[dim] = true]) Il semble que cette fonctionnalité ne soit plus utilisée. Nous prévoyons donc de la supprimer à l'avenir. (#1181).
• dot_product = dot_general(reversed_lhs_window, rhs, lhs_batching_dimensions=[], lhs_contracting_dimensions=input_spatial_dimensions + [input_feature_dimension], rhs_batching_dimensions=[], rhs_contracting_dimensions=kernel_spatial_dimensions + [kernel_input_feature_dimension]).

Si la valeur est feature_group_count > 1:

• lhses = split(lhs, feature_group_count, input_feature_dimension).
• rhses = split(rhs, feature_group_count, kernel_output_feature_dimension)
• results... = convolution(lhses..., rhses..., ..., feature_group_count=1, ...)
• result = concatenate(results, output_feature_dimension).

Si la valeur est batch_group_count > 1:

• lhses = split(lhs, batch_group_count, input_batch_dimension).
• rhses = split(rhs, batch_group_count, kernel_output_feature_dimension)
• results... = convolution(lhses..., rhses..., ..., batch_group_count=1, ...).
• result = concatenate(results, output_feature_dimension)

Pour les types quantifiés, exécute dequantize_op_quantize( lambda lhs, rhs: convolution(lhs, rhs, window_strides, padding, lhs_dilation, rhs_dilation, window_reversal, input_batch_dimension, input_feature_dimension, input_spatial_dimensions, kernel_input_feature_dimension, kernel_output_feature_dimension, kernel_spatial_dimensions, output_batch_dimension, output_feature_dimension, output_spatial_dimensions, feature_group_count, batch_group_count, precision_config), lhs, rhs, type(result)).

Pour les types hybrides quantifiés, exécute hybrid_dequantize_then_op( lambda lhs, rhs: convolution(lhs, rhs, window_strides, padding, lhs_dilation, rhs_dilation, window_reversal, input_batch_dimension, input_feature_dimension, input_spatial_dimensions, kernel_input_feature_dimension, kernel_output_feature_dimension, kernel_spatial_dimensions, output_batch_dimension, output_feature_dimension, output_spatial_dimensions, feature_group_count, batch_group_count, precision_config), lhs, rhs).

Entrées

Sorties

Contraintes

• (C1) N = rank(lhs) = rank(rhs).
• (C2) size(window_strides) = N - 2.
• (C3) 0 < window_strides.
• (C4) shape(padding) = [N - 2, 2].
• (C5) size(lhs_dilation) = N - 2.
• (C6) 0 < lhs_dilation.
• (C7) size(rhs_dilation) = N - 2.
• (C8) 0 < rhs_dilation.
• (C9) size(window_reversal) = N - 2.
• (C10) dim(lhs, input_batch_dimension) % batch_group_count = 0.
• (C11) dim(lhs, input_feature_dimension) % feature_group_count = 0.
• (C12) size(input_spatial_dimensions) = N - 2.
• (C13) Avec input_dimensions = [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension]: <ph type="x-smartling-placeholder">
  </ph>
  • is_unique(input_dimensions).
  • 0 <= input_dimensions < N.
• (C14) dim(rhs, kernel_input_feature_dimension) = dim(lhs, input_feature_dimension) / feature_group_count.
• (C15) dim(rhs, kernel_output_feature_dimension) % batch_group_count = 0.
• (C16) dim(rhs, kernel_output_feature_dimension) % feature_group_count = 0.
• (C17) size(kernel_spatial_dimensions) = N - 2.
• (C18) Avec kernel_dimensions = kernel_spatial_dimensions + [kernel_input_feature_dimension] + [kernel_output_feature_dimension]: <ph type="x-smartling-placeholder">
  </ph>
  • is_unique(kernel_dimensions).
  • 0 <= kernel_dimensions < N.
• (C19) size(output_spatial_dimensions) = N - 2.
• (C20) Donnée output_dimensions = [output_batch_dimension] + output_spatial_dimensions + [output_feature_dimension]: <ph type="x-smartling-placeholder">
  </ph>
  • is_unique(output_dimensions).
  • 0 <= output_dimensions < N.
• (C21) 0 < feature_group_count.
• (C22) 0 < batch_group_count.
• (C23) feature_group_count = 1 or batch_group_count = 1.
• (C24) size(precision_config) = 2.
• (C25) dim(result, result_dim) est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • dim(lhs, input_batch_dimension) / batch_group_count si result_dim = output_batch_dimension.
  • dim(rhs, kernel_output_feature_dimension) si result_dim = output_feature_dimension.
  • num_windows dans les autres cas, où:
  • output_spatial_dimensions[spatial_dim] = result_dim.
  • lhs_dim = input_spatial_dimensions[spatial_dim]
  • rhs_dim = kernel_spatial_dimensions[spatial_dim]
  • dilated_input_shape[lhs_dim] = dim(lhs, lhs_dim) = 0 ? 0 : (dim(lhs, lhs_dim) - 1) * lhs_dilation[spatial_dim] + 1
  • padded_input_shape[lhs_dim] = padding[spatial_dim, 0] + dilated_input_shape[lhs_dim] + padding[spatial_dim, 1]
  • dilated_window_shape[lhs_dim] = dim(rhs, rhs_dim) = 0 ? 0 : (dim(rhs, rhs_dim) - 1) * rhs_dilation[spatial_dim] + 1
  • is_empty_window[lhs_dim] = padded_input_shape[lhs_dim] = 0 || dilated_window_shape[lhs_dim] > padded_input_shape[lhs_dim]
  • num_windows = is_empty_window[lhs_dim] ? 0 : floor((padded_input_shape[lhs_dim] - dilated_window_shape[lhs_dim]) / window_strides[spatial_dim]) + 1.
• (C26) rank(result) = N.
• Si l'opération utilise des Tensors non quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C27) element_type(lhs) = element_type(rhs) = element_type(result).
• Si l'opération utilise des Tensors quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C28) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs).
  • (C29) Si is_per_axis_quantized(rhs), puis sur quantization_dimension(rhs) = kernel_output_feature_dimension.
  • (C30) Si is_per_axis_quantized(result), alors quantization_dimension(result) = output_feature_dimension
  • Si la valeur est is_quantized(lhs):
  • (C31) storage_type(lhs) = storage_type(rhs).
  • (C32) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C33) Si is_per_tensor_quantized(rhs), alors is_per_tensor_quantized(result)
  • Si la valeur est !is_quantized(lhs):
  • (C34) element_type(lhs) = expressed_type(rhs) = element_type(result).

Exemples

// %lhs: [[
//        [
//          [1], [2], [5], [6]
//        ],
//        [
//          [3], [4], [7], [8]
//        ],
//        [
//          [10], [11], [14], [15]
//        ],
//        [
//          [12], [13], [16], [17]
//        ]
//      ]]
//
// %rhs: [
//        [[[1]], [[1]], [[1]]],
//        [[[1]], [[1]], [[1]]],
//        [[[1]], [[1]], [[1]]]
//       ]
%result = "stablehlo.convolution"(%lhs, %rhs) {
  window_strides = array<i64: 4, 4>,
  padding = dense<0> : tensor<2x2xi64>,
  lhs_dilation = array<i64: 2, 2>,
  rhs_dilation = array<i64: 1, 1>,
  window_reversal = array<i1: false, false>,
  // In the StableHLO dialect, dimension numbers are encoded via:
  // `[<input dimensions>]x[<kernel dimensions>]->[output dimensions]`.
  // "b" is batch dimension, "f" is feature dimension,
  // "i" is input feature dimension, "o" is output feature dimension,
  // "0/1/etc" are spatial dimensions.
  dimension_numbers = #stablehlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  batch_group_count = 1 : i64,
  feature_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi64>, tensor<3x3x1x1xi64>) -> tensor<1x2x2x1xi64>
// %result: [[
//            [[10], [26]],
//            [[46], [62]]
//          ]]

Autres exemples

cosinus

Sémantique

Effectue une opération cosinus par élément sur le Tensor operand et génère une Tensor result. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les nombres à virgule flottante: cos d'IEEE-754.
• Pour les nombres complexes: cosinus complexe.
• Pour les types quantifiés: dequantize_op_quantize(cosine, operand, type(result)).

Entrées

Sorties

Contraintes

• (C1) baseline_type(operand) = baseline_type(result).

Exemples

// %operand: [
//            [0.0, 1.57079632],       // [0, pi/2]
//            [3.14159265, 4.71238898] // [pi, 3pi/2]
//           ]
%result = "stablehlo.cosine"(%operand) : (tensor<2x2xf32>) -> tensor<2x2xf32>
// %result: [[1.0, 0.0], [-1.0, 0.0]]

Autres exemples

count_leading_zeros

Sémantique

Effectue un comptage par élément du nombre de bits zéro au début dans operand. et produit un Tensor result.

Entrées

Sorties

Contraintes

• (C1) type(operand) = type(result).

Exemples

// %operand: [[0, 1], [128, -1]]
%result = "stablehlo.count_leading_zeros"(%operand) : (tensor<2x2xi64>) -> tensor<2x2xi64>
// %result: [[64, 63], [56, 0]]

Autres exemples

custom_call

Sémantique

Encapsule une opération call_target_name définie par l'implémentation qui accepte inputs et called_computations, et génère results. has_side_effect, backend_config et api_version peuvent être utilisés pour fournir définies par l'implémentation.

Pour le moment, cette opération contient une collection assez désorganisée de de métadonnées qui reflètent l'évolution organique de son fonctionnement équivalent dans le compilateur XLA. À l'avenir, nous prévoyons d'unifier ces métadonnées (#741)

Entrées

Sorties

Exemples

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

diviser

Sémantique

Effectue une division par élément des Tensors lhs et diviseur rhs. et génère un Tensor result. En fonction du type d'élément, effectue les opérations suivantes:

• Pour les entiers: division des entiers, qui produit le quotient algébrique avec n'importe lequel partie fractionnaire supprimée.
• Pour les nombres à virgule flottante: division d'IEEE-754.
• Pour les nombres complexes: division complexe.
• Pour les types quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • dequantize_op_quantize(divide, lhs, rhs, type(result)).

Entrées

Sorties

Contraintes

• (C1) baseline_type(lhs) = baseline_type(rhs) = baseline_type(result).

Exemples

// %lhs: [17.1, -17.1, 17.1, -17.1]
// %rhs: [3.0, 3.0, -3.0, -3.0]
%result = "stablehlo.divide"(%lhs, %rhs) : (tensor<4xf32>, tensor<4xf32>) -> tensor<4xf32>
// %result: [5.66666651, -5.66666651, -5.66666651, 5.66666651]

Autres exemples

dot_general

Sémantique

Calcule les produits scalaires entre les tranches de lhs et les tranches de rhs, et génère une Tensor result.

Plus formellement, result[result_index] = dot_product, où:

• lhs_result_dimensions = [d for d in axes(lhs) and d not in lhs_batching_dimensions and d not in lhs_contracting_dimensions].
• rhs_result_dimensions = [d for d in axes(rhs) and d not in rhs_batching_dimensions and d not in rhs_contracting_dimensions].
• result_batching_index + result_lhs_index + result_rhs_index = result_index où size(result_batching_index) = size(lhs_batching_dimensions), size(result_lhs_index) = size(lhs_result_dimensions) et size(result_rhs_index) = size(rhs_result_dimensions)
• transposed_lhs = transpose(lhs, lhs_batching_dimensions + lhs_result_dimensions + lhs_contracting_dimensions).
• transposed_lhs_slice = slice(transposed_lhs, result_batching_index + result_lhs_index + [:, ..., :])
• reshaped_lhs_slice = reshape(transposed_lhs_slice, dims(lhs, lhs_contracting_dimensions))
• transposed_rhs = transpose(rhs, rhs_batching_dimensions + rhs_result_dimensions + rhs_contracting_dimensions)
• transposed_rhs_slice = slice(transposed_rhs, result_batching_index + result_rhs_index + [:, ..., :])
• reshaped_rhs_slice = reshape(transposed_rhs_slice, dims(rhs, rhs_contracting_dimensions)).
• dot_product = reduce( inputs=[multiply(reshaped_lhs_slice, reshaped_rhs_slice)], init_values=[constant(0, element_type(result))], dimensions=range(size(lhs_contracting_dimensions)), body=lambda x, y: add(x, y))

Pour les types quantifiés, exécute dequantize_op_quantize( lambda lhs, rhs: dot_general(lhs, rhs, lhs_batching_dimensions, rhs_batching_dimensions, lhs_contracting_dimensions, rhs_contracting_dimensions, precision_config), lhs, rhs, type(result)).

Pour les types hybrides quantifiés, exécute hybrid_dequantize_then_op( lambda lhs, rhs: dot_general(lhs, rhs, lhs_batching_dimensions, rhs_batching_dimensions, lhs_contracting_dimensions, rhs_contracting_dimensions, precision_config), lhs, rhs).

precision_config contrôle le compromis entre vitesse et précision pour sur les backends d'accélérateur. Il peut s'agir de l'un des éléments suivants (au niveau Dans un premier temps, la sémantique de ces valeurs d'énumération est sous-spécifiée, mais nous n'avons prévoient de résoudre ce problème dans #755):

• DEFAULT: calcul le plus rapide, mais approximation la moins précise du numéro d'origine.
• HIGH: calcul plus lent, mais approximation plus précise du numéro d'origine.
• HIGHEST: calcul le plus lent, mais approximation la plus précise du numéro d'origine.

Un DotAlgorithm définit les principales propriétés de l'algorithme utilisé pour implémenter l'opération par point, qui définit également la précision. Si l'attribut de l'algorithme sont définis, l'élément precision_config doit être défini sur DEFAULT. DotAlgorithms n'ont pas de valeur par défaut, car les paramètres par défaut sont définis. Par conséquent, tous les champs de l'algorithme peuvent être définis sur None pour spécifier une algorithme vide avec points, qui utilisera plutôt la valeur precision_config.

Les champs DotAlgorithm incluent:

• lhs_precision_type et rhs_precision_type, la précision des fonctions LHS et à droite de l'opération sont arrondies. Les types de précision sont indépendants des de stockage des entrées et des sorties.
• accumulation_type est la précision utilisée pour l'accumulation.
• lhs_component_count, rhs_component_count et num_primitive_operations s'appliquent lorsque nous faisons un algorithme qui décompose le LHS et/ou le RHS en plusieurs composants et effectue plusieurs tâches sur ces valeurs : généralement pour émuler une précision plus élevée (par exemple, Exploiter le type de données bfloat16 de l'intelligence artificielle pour effectuer des calculs de plus grande précision: bf16_6x, tf32_3x, etc.). Pour les algorithmes sans décomposition, ces valeurs doit être défini sur 1.
• allow_imprecise_accumulation pour spécifier si l'accumulation dans une précision inférieure est autorisé pour certaines étapes (par exemple, CUBLASLT_MATMUL_DESC_FAST_ACCUM).

Exemples d'attributs DotAlgorithm:

// Inputs are casted to tf32, and then accumulated in f32:
{lhs_precision_type = tf32,
 rhs_precision_type = tf32,
 accumulation_type = f32,
 lhs_component_count = 1,
 rhs_component_count = 1,
 num_primitive_operations = 1,
 allow_imprecise_accumulation = false}


// bf16_6x: each input is decomposed to 3 bf16 components, then 6 dot operations are done on those components, and the result is accumulated in f32.
{lhs_precision_type = bf16,
 rhs_precision_type = bf16,
 accumulation_type = f32,
 lhs_component_count = 3,
 rhs_component_count = 3,
 num_primitive_operations = 6,
 allow_imprecise_accumulation = false}


// Inputs are (casted to) f8e5m2, and we accumulate in f32, but for some steps we may accumulate in lower precision.
{lhs_precision_type = f8e5m2,
 rhs_precision_type = f8e5m2,
 accumulation_type = f32,
 lhs_component_count = 1,
 rhs_component_count = 1,
 num_primitive_operations = 1,
 allow_imprecise_accumulation = true}

C'est aux implémentations de décider quelles combinaisons sont acceptées. Dans général, il n'est pas garanti que chaque algorithme soit compatible type d'accélérateur par le consommateur de StableHLO. Si un algorithme donné n'est pas une erreur doit être générée, plutôt que de revenir à une alternative. La vérification StableHLO s'efforcera au mieux, empêchant ainsi les algorithmes qui ne sont pas compatibles avec aucun matériel.

Voir xla_data.proto > Algorithm pour certaines valeurs d'algorithme acceptées. Le ticket n° 2483 capture le plan pour créer un document centralisé sur les algorithmes pris en charge par le backend.

Entrées

Sorties

Contraintes

• (C1) size(lhs_batching_dimensions) = size(rhs_batching_dimensions).
• (C2) size(lhs_contracting_dimensions) = size(rhs_contracting_dimensions).
• (C3) is_unique(lhs_batching_dimensions + lhs_contracting_dimensions).
• (C4) is_unique(rhs_batching_dimensions + rhs_contracting_dimensions).
• (C5) 0 <= lhs_batching_dimensions < rank(lhs).
• (C6) 0 <= lhs_contracting_dimensions < rank(lhs).
• (C7) 0 <= rhs_batching_dimensions < rank(rhs).
• (C8) 0 <= rhs_contracting_dimensions < rank(rhs).
• (C9) dim(lhs, lhs_batching_dimensions...) = dim(rhs, rhs_batching_dimensions...).
• (C10) dim(lhs, lhs_contracting_dimensions...) = dim(rhs, rhs_contracting_dimensions...).
• (C11) size(precision_config) = 2.
• (C12) shape(result) = dim(lhs, lhs_batching_dimensions) + dim(lhs, lhs_result_dimensions) + dim(rhs, rhs_result_dimensions).
• Si l'opération utilise des Tensors non quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C13) element_type(lhs) = element_type(rhs).
• Si l'opération utilise des Tensors quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C14) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs).
  • (C15) zero_points(rhs) = 0.
  • (C16) Si is_per_axis_quantized(rhs), alors quantization_dimension(rhs) pas dans rhs_contracting_dimensions.
  • Si la valeur est is_quantized(lhs):
  • (C17) storage_type(lhs) = storage_type(rhs).
  • (C18) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C19) Si is_per_tensor_quantized(rhs), alors is_per_tensor_quantized(result)
  • Si la valeur est !is_quantized(lhs):
  • (C20) element_type(lhs) = expressed_type(rhs) = element_type(result).
• Si la valeur est !is_empty_algorithm(lhs_precision_type, rhs_precision_type, accumulation_type, lhs_component_count, rhs_component_count, num_primitive_operations allow_imprecise_accumulation): <ph type="x-smartling-placeholder">
  </ph>
  • (C21) precision_config... = DEFAULT.
  • (C22) 0 < lhs_component_count.
  • (C23) 0 < rhs_component_count.
  • (C24) 0 < num_primitive_operations.

Exemples

// %lhs: [
//        [[1, 2],
//         [3, 4]],
//        [[5, 6],
//         [7, 8]]
//       ]
// %rhs: [
//        [[1, 0],
//         [0, 1]],
//        [[1, 0],
//         [0, 1]]
//       ]
%result = "stablehlo.dot_general"(%lhs, %rhs) {
  dot_dimension_numbers = #stablehlo.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>],
  algorithm = #stablehlo.dot_algorithm<
    lhs_precision_type = tf32,
    rhs_precision_type = tf32,
    accumulation_type = f32,
    lhs_component_count = 1,
    rhs_component_count = 1,
    num_primitive_operations = 1,
    allow_imprecise_accumulation = false
  >
} : (tensor<2x2x2xi64>, tensor<2x2x2xi64>) -> tensor<2x2x2xi64>
// %result: [
//           [[1, 2],
//            [3, 4]],
//           [[5, 6],
//            [7, 8]]
//          ]

Autres exemples

dynamic_broadcast_in_dim

Sémantique

Cette opération est fonctionnellement identique à broadcast_in_dim op, mais la forme du résultat est spécifiée dynamiquement via output_dimensions.

L'opération accepte également les attributs facultatifs known_expanding_dimensions et known_non_expanding_dimensions. pour exprimer des connaissances statiques sur le comportement d'expansion des dimensions. Si aucune valeur n'est spécifiée, toutes les dimensions sont supposées pouvoir s'étendre.

Entrées

Sorties

Contraintes

• (C1) La valeur element_type(result) est donnée par: <ph type="x-smartling-placeholder">
  </ph>
  • element_type(operand), si !is_per_axis_quantized(operand).
  • element_type(operand), à l'exception de quantization_dimension(operand), scales(operand) et zero_points(operand) peuvent être différents de quantization_dimension(result), scales(result) et zero_points(result) resp., sinon.
• (C2) size(broadcast_dimensions) = rank(operand).
• (C3) 0 <= broadcast_dimensions < rank(result).
• (C4) is_unique(broadcast_dimensions).
• (C5) Pour tous les d de axes(operand): <ph type="x-smartling-placeholder">
  </ph>
  • dim(operand, d) = 1 ou
  • dim(operand, d) = dim(result, broadcast_dimensions[d]).
• (C6) Si is_per_axis_quantized(result): <ph type="x-smartling-placeholder">
  </ph>
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)].
  • Si la valeur est dim(operand, quantization_dimension(operand)) = 1, alors scales(result)[i] = scales(operand)[0] and zero_points(result)[i] = zero_points(operand)[0] for i in range(dim(result, quantization_dimension(result)))
• (C7) size(output_dimensions) = rank(result).
• (C8) is_unique(known_expanding_dimensions + known_non_expanding_dimensions).
• (C9) 0 <= known_expanding_dimensions < rank(operand).
• (C10) 0 <= known_non_expanding_dimensions < rank(operand).

Exemples

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

Autres exemples

dynamic_conv

Sémantique

Cette opération est fonctionnellement identique à Convolution op, mais la marge intérieure est spécifiée dynamiquement via padding.

Entrées

Sorties

Contraintes

• (C1) N = rank(lhs) = rank(rhs).
• (C2) size(window_strides) = N - 2.
• (C3) 0 < window_strides.
• (C4) shape(padding) = [N - 2, 2].
• (C5) size(lhs_dilation) = N - 2.
• (C6) 0 < lhs_dilation.
• (C7) size(rhs_dilation) = N - 2.
• (C8) 0 < rhs_dilation.
• (C9) size(window_reversal) = N - 2.
• (C10) dim(lhs, input_batch_dimension) % batch_group_count = 0.
• (C11) dim(lhs, input_feature_dimension) % feature_group_count = 0.
• (C12) size(input_spatial_dimensions) = N - 2.
• (C13) Avec input_dimensions = [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension]: <ph type="x-smartling-placeholder">
  </ph>
  • is_unique(input_dimensions).
  • 0 <= input_dimensions < N.
• (C14) dim(rhs, kernel_input_feature_dimension) = dim(lhs, input_feature_dimension) / feature_group_count.
• (C15) dim(rhs, kernel_output_feature_dimension) % batch_group_count = 0.
• (C16) dim(rhs, kernel_output_feature_dimension) % feature_group_count = 0.
• (C17) size(kernel_spatial_dimensions) = N - 2.
• (C18) Avec kernel_dimensions = kernel_spatial_dimensions + [kernel_input_feature_dimension] + [kernel_output_feature_dimension]: <ph type="x-smartling-placeholder">
  </ph>
  • is_unique(kernel_dimensions).
  • 0 <= kernel_dimensions < N.
• (C19) size(output_spatial_dimensions) = N - 2.
• (C20) Donnée output_dimensions = [output_batch_dimension] + output_spatial_dimensions + [output_feature_dimension]: <ph type="x-smartling-placeholder">
  </ph>
  • is_unique(output_dimensions).
  • 0 <= output_dimensions < N.
• (C21) 0 < feature_group_count.
• (C22) 0 < batch_group_count.
• (C23) feature_group_count = 1 or batch_group_count = 1.
• (C24) size(precision_config) = 2.
• (C25) dim(result, result_dim) est défini comme suit: <ph type="x-smartling-placeholder">
  </ph>
  • dim(lhs, input_batch_dimension) / batch_group_count si result_dim = output_batch_dimension.
  • dim(rhs, kernel_output_feature_dimension) si result_dim = output_feature_dimension.
  • num_windows dans les autres cas, où:
  • output_spatial_dimensions[spatial_dim] = result_dim.
  • lhs_dim = input_spatial_dimensions[spatial_dim]
  • rhs_dim = kernel_spatial_dimensions[spatial_dim]
  • dilated_input_shape[lhs_dim] = dim(lhs, lhs_dim) = 0 ? 0 : (dim(lhs, lhs_dim) - 1) * lhs_dilation[spatial_dim] + 1
  • padded_input_shape[lhs_dim] = padding[spatial_dim, 0] + dilated_input_shape[lhs_dim] + padding[spatial_dim, 1]
  • dilated_window_shape[lhs_dim] = dim(rhs, rhs_dim) = 0 ? 0 : (dim(rhs, rhs_dim) - 1) * rhs_dilation[spatial_dim] + 1
  • is_empty_window[lhs_dim] = padded_input_shape[lhs_dim] = 0 || dilated_window_shape[lhs_dim] > padded_input_shape[lhs_dim]
  • num_windows = is_empty_window[lhs_dim] ? 0 : floor((padded_input_shape[lhs_dim] - dilated_window_shape[lhs_dim]) / window_strides[spatial_dim]) + 1.
• (C26) rank(result) = N.
• Si l'opération utilise des Tensors non quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C27) element_type(lhs) = element_type(rhs) = element_type(result).
• Si l'opération utilise des Tensors quantifiés: <ph type="x-smartling-placeholder">
  </ph>
  • (C28) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs).
  • (C29) Si is_per_axis_quantized(rhs), puis sur quantization_dimension(rhs) = kernel_output_feature_dimension.
  • (C30) Si is_per_axis_quantized(result), alors quantization_dimension(result) = output_feature_dimension
  • Si la valeur est is_quantized(lhs):
  • (C31) storage_type(lhs) = storage_type(rhs).
  • (C32) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C33) Si is_per_tensor_quantized(rhs), alors is_per_tensor_quantized(result)
  • Si la valeur est !is_quantized(lhs):
  • (C34) element_type(lhs) = expressed_type(rhs) = element_type(result).

Exemples

// %lhs: [[
//        [[1], [2], [5], [6]],
//        [[3], [4], [7], [8]],
//        [[10], [11], [14], [15]],
//        [[12], [13], [16], [17]]
//      ]]
//
// %rhs: [
//         [[[1]], [[1]], [[1]]],
//         [[[1]], [[1]], [[1]]],
//         [[[1]], [[1]], [[1]]]
//        ]
// %padding: [[1, 1],
//            [1, 1]]
%result = "stablehlo.dynamic_conv"(%lhs, %rhs, %padding) {
  window_strides = array<i64: 4, 4>,
  lhs_dilation = array<i64: 2, 2>,
  rhs_dilation = array<i64: 1, 1>,
  window_reversal = array<i1: false, false>,
  dimension_numbers = #stablehlo.conv<raw
    input_batch_dimension = 0,
    input_feature_dimension = 3,
    input_spatial_dimensions = [0, 1],
    kernel_input_feature_dimension = 2,
    kernel_output_feature_dimension = 3,
    kernel_spatial_dimensions = [0, 1],
    output_batch_dimension = 0,
    output_feature_dimension = 3,
    output_spatial_dimensions = [1, 2]
  >,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi64>, tensor<3x3x1x1xi64>, tensor<2x2xi64>) -> tensor<1x2x2x1xi64>
// %result: [[
//            [[1], [5]],
//            [[10], [14]]
//          ]]