StableHLO-Spezifikation

StableHLO ist ein Operation-Set für High-Level-Operationen (HLO) in Maschinen, Lernmodelle (ML). StableHLO dient als Portabilitätsebene zwischen verschiedenen ML-Frameworks und ML-Compiler: ML-Frameworks, die StableHLO-Programme erstellen sind mit ML-Compilern kompatibel, die StableHLO-Programme nutzen.

Unser Ziel ist es, die ML-Entwicklung zu vereinfachen und zu beschleunigen, Interoperabilität zwischen verschiedenen ML-Frameworks (wie TensorFlow, JAX und PyTorch) und ML-Compiler (wie XLA und IREE). Gegen dieses Ende enthält eine Spezifikation für die Programmiersprache StableHLO.

Diese Spezifikation enthält drei Hauptabschnitte. Zunächst sollte der Im Abschnitt Programme wird die Struktur von StableHLO-Programmen beschrieben. die aus StableHLO-Funktionen bestehen, die wiederum aus StableHLO-Operationen bestehen. Innerhalb dieser Struktur gibt der Abschnitt Ops die Semantik der für einzelne Vorgänge. Der Abschnitt Ausführung enthält Semantik für alle die diese Vorgänge gemeinsam in einem Programm ausführen. Die Funktion Im Abschnitt Notation wird die im gesamten Dokument verwendete Notation erläutert. Spezifikation zu ändern.

Um die Spezifikation einer früheren StableHLO-Version anzuzeigen, öffnen Sie das Repository im getaggte Freigabe von Interesse. Beispiel: StableHLO v0.19.0 Spec Informationen zum Anzeigen von Änderungen, die bei jedem Nebenversions-Bump von StableHLO vorgenommen wurden, finden Sie unter das Versionsprotokoll in VhloDialect.td

Programme

Program ::= {Func}

StableHLO-Programme bestehen aus einer beliebigen Anzahl von StableHLO-Funktionen. Unten sehen Sie ein Beispielprogramm mit der Funktion @main mit 3 Eingaben (%image, %weights und %bias) und 1 Ausgabe. Der Hauptteil der Funktion hat 6 Vorgänge.

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>) -> ()
}

Funktionen

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

StableHLO-Funktionen, die auch als benannte Funktionen bezeichnet werden, haben eine Kennung, Ein-/Ausgaben und einen Textkörper. Für die Zukunft planen wir, Es werden zusätzliche Metadaten für Funktionen eingeführt, um eine bessere Kompatibilität zu erreichen. mit HLO (#425, #626 Nr. 740 #744).

IDs

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

StableHLO-Kennungen ähneln Kennungen in vielen Programmiersprachen. Sprachen, mit zwei Besonderheiten: 1) Alle Kennungen haben Sieigzeichen, die verschiedene Arten von Kennungen unterscheiden, 2) Wertkennungen vollständig numerisch, um die Generierung von StableHLO-Programmen zu vereinfachen.

Typen

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

StableHLO-Typen werden in Werttypen kategorisiert, die auch als First-class-Typen), die StableHLO-Werte und Nicht-Werttypen darstellen. die andere Programmelemente beschreiben. StableHLO-Typen ähneln Typen in viele Programmiersprachen, wobei die größte Besonderheit die Funktionen von StableHLO sind. domainspezifischen Charakter, der zu ungewöhnlichen Ergebnissen führt (z.B. skalare Typen) sind keine Werttypen).

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

Tensor-Typen stellen Tensoren dar, also mehrdimensionale Arrays. Sie haben eine Form und einen Unterelementtyp, wobei eine Form für nicht negative oder unbekannte Dimensionsgrößen in aufsteigender Reihenfolge der zugehörigen Dimensionen (auch Achsen genannt) von 0 bis R-1 nummeriert. Die Anzahl der Dimensionen R wird als Rang bezeichnet. Beispiel: tensor<2x3xf32> ist einen Tensortyp mit der Form 2x3 und dem Elementtyp f32. Es hat zwei Dimensionen. (oder, mit anderen Worten, zwei Achsen) – die 0. Dimension und die erste Dimension, deren Größen 2 und 3 sind. Sein Rang ist 2.

Formen können teilweise oder vollständig unbekannt (dynamisch) sein, z.B. tensor<?x2xf64> ist teilweise unbekannt und tensor<?x?xf64> ist vollständig unbekannt. Dynamisch Dimensionsgrößen durch ? dargestellt. Die Rangfolge von Formen kann nicht aufgehoben werden.

Zukünftig möchten wir die Tensortypen Größen und Elementtypen beispielsweise, um Layouts (#629) und Datendichte (#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

Quantisierte Elementtypen stehen für ganzzahlige Werte eines Speichertyps in den Bereich von storage_min bis storage_max (einschließlich), die sich auf Gleitkommawerte eines ausgedrückten Typs. Für einen gegebenen ganzzahligen Wert i kann der entsprechende Gleitkommawert f als f = (i - zero_point) * scale, wobei scale und zero_point aufgerufen werden Quantisierungsparametern. storage_min und storage_max sind optional in der Grammatik, aber haben die Standardwerte min_value(storage_type) und max_value(storage_type). Quantisierte Elementtypen haben die folgenden Einschränkungen:

• (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) Wenn is_empty(quantization_dimension), dann size(scales) = 1.
• (C11) 0 <= quantization_dimension.

Derzeit ist QuantizationScale eine Gleitkommakonstante, es gibt jedoch starkes Interesse an ganzzahlbasierten Skalen, dargestellt durch Multiplikatoren und Schichten. Wir möchten dies in naher Zukunft untersuchen (#1404)

Die Semantik von QuantizationZeroPoint wird derzeit diskutiert. einschließlich Typ, Werten und ob es nur eine oder in einem quantisierten Tensortyp. Basierend auf den zum Ergebnis dieser Diskussion kann sich die Spezifikation um Nullpunkte in der Zukunft erfolgen (#1405).

Eine weitere laufende Diskussion zur Semantik von QuantizationStorageMin und QuantizationStorageMax, um zu bestimmen, ob Einschränkungen auf diese Werte und die Werte von quantisierten Tensoren angewendet. (#1406)

Schließlich wollen wir erkunden, wie unbekannte Skalen und Nullen Ähnlich wie die Darstellung von unbekannten Dimensionsgrößen (#1407)

Quantisierte Tensortypen stellen Tensoren mit quantisierten Elementen dar. Diese Tensoren sind genau die gleichen wie reguläre Tensoren, mit der Ausnahme, dass ihre Elemente quantisierte Elementtypen anstelle von regulären Elementtypen haben.

In quantisierten Tensoren kann die Quantisierung per-Tensor erfolgen, was bedeutet, Ein scale und zero_point für den gesamten Tensor oder kann pro Achse sein, Dies bedeutet, dass bei mehreren scales und zero_points ein Paar pro Segment von um eine bestimmte Dimension quantization_dimension zu erhalten. Formals in einem Tensor-t Mit der Quantisierung pro Achse gibt es dim(t, quantization_dimension) Slices der quantization_dimension: t[:, ..., 0, ..., :], t[:, ..., 1, ..., :], Alle Elemente im i. Segment verwenden scales[i] und zero_points[i] als ihre Quantisierungsparameter. Quantisierte Tensortypen haben Folgendes: Einschränkungen:

• Für die Quantisierung pro Tensor: <ph type="x-smartling-placeholder">
  </ph>
  • Keine zusätzlichen Einschränkungen.
• Für die Quantisierung pro Achse: <ph type="x-smartling-placeholder">
  </ph>
  • (C12) quantization_dimension < rank(self).
  • (C13) dim(self, quantization_dimension) = size(scales).

TokenType ::= 'token'

Tokentypen stellen Tokens dar, d. h. intransparente Werte, die erzeugt und verbraucht werden. durch einige Operationen. Tokens werden zum Festlegen der Ausführungsreihenfolge für Vorgänge verwendet wie im Abschnitt Ausführung beschrieben.

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

Tupeltypen stehen für Tupel, also heterogene Listen. Tupel sind ein Vermächtnis die nur aus Gründen der Kompatibilität mit HLO besteht. In HLO werden Tupel zur Darstellung variadischer Ein- und Ausgaben. In StableHLO werden variadische Eingaben und Ausgaben werden nativ unterstützt und die einzige Verwendung von Tupeln in StableHLO besteht darin, umfassende Darstellung von HLO ABI, wo z.B. T, tuple<T> und tuple<tuple<T>> kann je nach Implementierung. Wir planen, in Zukunft Änderungen am HLO ABI vorzunehmen. mit der wir Tupeltypen aus StableHLO entfernen können, (#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'

Elementtypen stellen Elemente von Tensortypen dar. Anders als bei vielen Programmen Sprachen sind, sind diese Typen in StableHLO nicht an erster Stelle. Das bedeutet, dass StableHLO-Programme können Werte dieser Typen nicht direkt darstellen (daher Es ist idiomatisch, Skalarwerte des Typs T mit einem 0-dimensionalen Tensor darzustellen. Werte des Typs tensor<T>) enthält.

• Der Boolesche Typ steht für die booleschen Werte true und false.
• Ganzzahltypen können entweder vorzeichen (si) oder unvorzeichent (ui) sein und haben eine der unterstützten Bitbreiten (2, 4, 8, 16, 32 oder 64). Vorzeichenbehaftete siN-Typen stellen ganzzahlige Werte von -2^(N-1) bis 2^(N-1)-1 dar einschließlich und vorzeichenlose uiN-Typen stellen Ganzzahlwerte von 0 bis 2^N-1 einschließlich.
• Es gibt folgende Gleitkommatypen: <ph type="x-smartling-placeholder">
  </ph>
  • f8E4M3FN- und f8E5M2-Typen, die den entsprechenden E4M3- und E5M2-Codierungen des FP8-Formats beschrieben in FP8-Formate für Deep Learning.
  • Typen f8E4M3FNUZ und f8E5M2FNUZ, die E4M3 und E5M2 entsprechen Codierungen der FP8-Formate, die in Numerische 8-Bit-Formate für neuronale Deep-Learning-Netzwerke.
  • Typ f8E4M3B11FNUZ, der der E4M3-Codierung der FP8-Formate entspricht beschrieben in Training und Inferenz für neuronale Deep-Learning-Netzwerke (Hybrid 8-Bit Floating Point, HFP8).
  • bf16-Typ, der dem bfloat16-Format entspricht, das beschrieben wird in BFloat16: Das Geheimnis für hohe Leistung auf Cloud TPUs
  • f16-, f32- und f64-Typen, die dem jeweiligen Typ entsprechen binary16 ("halbe Genauigkeit"), binary32 ("einfache Genauigkeit") und binary64-Formate („Double Precision“), die unter IEEE 754-Standard.
  • Der Typ tf32 entspricht dem TensorFloat32-Format. und wird in StableHLO nur eingeschränkt unterstützt.
• Komplexe Typen stehen für komplexe Werte, die einen reellen Teil haben und einem imaginären Teil desselben Elementtyps. Unterstützt, komplex Typen sind complex<f32> (beide Teile sind vom Typ f32) und complex<f64> (beide Teile sind vom Typ f64).

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

Funktionstypen stellen sowohl benannte als auch anonyme Funktionen dar. Sie haben Eingabetypen (die Liste der Typen auf der linken Seite von ->) und Ausgabetypen (die Liste der Typen rechts von ->). In vielen Programmierungen sind Funktionstypen erstklassig, aber in StableHLO nicht.

StringType ::= 'string'

Der Stringtyp steht für Bytesequenzen. Anders als bei vielen Programmen ist der Zeichenfolgentyp in StableHLO nicht die erste Klasse und wird nur verwendet, geben statische Metadaten für Programmelemente an.

Vorgänge

StableHLO-Operationen (auch Ops genannt) stellen eine geschlossene Menge dar. übergeordneten Abläufen in ML-Modellen. Wie bereits erwähnt, StableHLO-Syntax ist stark von MLIR inspiriert, was nicht unbedingt die am stärksten eine ergonomische Alternative, ist aber wohl am besten für das Ziel von StableHLO geeignet, und mehr Interoperabilität zwischen ML-Frameworks und ML-Compilern schaffen.

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

StableHLO-Operationen (auch Operations genannt) haben einen Namen, Ein-/Ausgaben und eine Signatur. Der Name besteht aus dem Präfix stablehlo. und Eine mnemonische Kennzeichnung, die einen der unterstützten Vorgänge eindeutig identifiziert. Siehe unten für eine umfassende Liste aller unterstützten Operationen.

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

Betriebsabläufe verbrauchen Eingaben und erzeugen Ausgaben. Eingaben sind kategorisiert in Eingabewerte (berechnet während der Ausführung), Eingabefunktionen (bereitgestellt statisch, da StableHLO-Funktionen keine First-Class-Werte sind) und Eingabeattribute (auch statisch bereitgestellt). Die Art der Ein- und Ausgaben von einem Opaion konsumiert und produziert wird, hängt von seiner Eselsbrücke ab. Beispiel: add op verbraucht 2 Eingabewerte und erzeugt 1 Ausgabewert. Im Vergleich dazu Der Vorgang select_and_scatter verbraucht 3 Eingabewerte, 2 Eingabefunktionen und 3 Eingabeattribute.

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

Eingabefunktionen (auch anonyme Funktionen genannt) sind sehr ähnlich wie benannte Funktionen, mit dem Unterschied, dass sie keine Kennung haben (daher der Name „anonymous“), 2) es werden keine Ausgabetypen deklariert (Ausgabetypen sind aus dem return-Vorgang in der Funktion abgeleitet wird).

Die Syntax für Eingabefunktionen enthält einen derzeit nicht verwendeten Teil (siehe Unused-Produktion oben), das aus Gründen der MLIR-Kompatibilität gedacht ist. In MLIR gibt es ein allgemeineres Konzept von „Regionen“, die mehrere „Blöcke“ haben kann, der Operationen sind über Jump Ops miteinander verbunden. Die IDs dieser Blöcke entsprechen an die Unused-Produktion, damit sie voneinander unterschieden werden können. StableHLO hat keine Jump Ops, sodass der entsprechende Teil der MLIR-Syntax nicht verwendet werden (aber noch vorhanden ist).

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

Eingabeattribute haben einen Namen und einen Wert. Dies ist einer der unterstützten Werte. Konstanten. Sie sind die primäre Möglichkeit, statische Metadaten für das Programm anzugeben. Elemente. Beim Vorgang concatenate wird beispielsweise das Attribut dimension verwendet, gibt die Dimension an, mit der die Eingabewerte verkettet werden. In ähnlicher Weise Der Vorgang slice verwendet mehrere Attribute wie start_indices und limit_indices , um die Grenzen anzugeben, die zum Aufteilen des Eingabewerts verwendet werden.

Derzeit enthalten StableHLO-Programme, die derzeit frei sind, manchmal Attribute die in diesem Dokument nicht beschrieben sind. Für die Zukunft planen wir, diese Attribute entweder in das StableHLO-Opset aufnehmen oder sie StableHLO-Programmen. In der Zwischenzeit finden Sie hier eine Liste mit diesen Attribute:

• layout (#629)
• mhlo.frontend_attributes (#628)
• mhlo.sharding (#619)
• output_operand_aliases (#740)
• Standortmetadaten (#594)

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

Die Vorgangssignatur besteht aus den Typen aller Eingabewerte (der Liste der Typen der linken Seite von ->) und die Typen aller Ausgabewerte (die Liste der auf der rechten Seite von ->. Genau genommen sind Eingabetypen redundant sind und Ausgabetypen fast immer redundant sind. Bei den meisten StableHLO-Operationen lassen sich Ausgabetypen aus Eingaben ableiten). Trotzdem Signatur ist absichtlich Teil der StableHLO-Syntax für die Kompatibilität mit MLIR.

Unten sehen Sie ein Beispiel für eine Operation mit der mnemonischen Kennzeichnung select_and_scatter. Es verbraucht 3 Eingabewerte (%operand, %source und %init_value), 2 Eingabefunktionen und 3 Eingabeattribute (window_dimensions, window_strides und padding). Beachten Sie, dass die Signatur des Vorgangs nur die Typen der Eingabewerte enthält. Dies gilt jedoch nicht für die Art der Eingabefunktionen und Attribute, die inline bereitgestellt werden.

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

Konstanten

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

StableHLO-Konstanten haben ein Literal und einen Typ, die zusammen für einen StableHLO-Wert. Im Allgemeinen ist der Typ Teil der konstanten Syntax, außer wenn sie eindeutig ist (z.B. hat eine boolesche Konstante eindeutig den Typ i1, während eine ganzzahlige Konstante mehrere mögliche Typen haben kann).

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

Boolesche Konstanten stellen die booleschen Werte true und false dar. Boolescher Wert haben Konstanten den Typ i1.

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

Ganzzahlkonstanten stellen Ganzzahlwerte über Strings dar, die Dezimalzahlen oder Hexadezimalschreibweise. Andere Basen, z.B. binär oder oktal, werden nicht unterstützt. Für Ganzzahlkonstanten gelten die folgenden Einschränkungen:

• (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]

Gleitkommakonstanten stellen Gleitkommawerte über Strings dar, die Dezimalschreibweise oder wissenschaftliche Schreibweise verwenden. Außerdem kann die hexadezimale Notation werden die zugrunde liegenden Bits direkt im Gleitkommaformat eines mit dem entsprechenden Typ aus. Für Gleitkommakonstanten gelten die folgenden Einschränkungen:

• (C1) Wenn die Hexadezimalschreibweise nicht verwendet wird, is_wellformed(float_literal, float_type)
• (C2) Bei Verwendung der Hexadezimalschreibweise size(hexadecimal_digits) = num_bits(float_type) / 4

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

Komplexe Konstanten stellen komplexe Werte mithilfe von Listen eines reellen Teils dar. (kommt an erster Stelle) und ein imaginärer Teil (an zweiter Stelle). Beispiel: (1.0, 0.0) : complex<f32> steht für 1.0 + 0.0i und (0.0, 1.0) : complex<f32> repräsentiert 0.0 + 1.0i. Die Reihenfolge, in der diese werden -Teile dann im Arbeitsspeicher gespeichert. Komplexe Konstanten haben folgende Einschränkungen:

• (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

Tensor-Konstanten stellen Tensor-Werte mithilfe verschachtelter Listen dar, die über NumPy-Notation. Beispiel: dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> stellt einen Tensorwert mit der folgenden Zuordnung von Indexen zu Elementen dar: {0, 0} => 1, {0, 1} => 2, {0, 2} => 3, {1, 0} => 4, {1, 1} => 5, {1, 2} => 6 Die Reihenfolge, in der diese Elemente dann im Arbeitsspeicher gespeichert werden, Implementierung definiert. Für Tensorkonstanten gelten die folgenden Einschränkungen:

• (C1) has_syntax(tensor_literal, element_type(tensor_type)), wobei gilt: <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)), wobei gilt: <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:]).
  • Andernfalls false.

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

Quantisierte Tensorkonstanten stellen quantisierte Tensorwerte mit denselben Werten Notation als Tensorkonstanten, wobei Elemente als Konstanten ihrer Speichertyp. Quantisierte Tensorkonstanten haben die folgenden Einschränkungen:

• (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))

Stringliterale bestehen aus Byte, die mithilfe von ASCII-Zeichen und Escapesequenzen verwenden. Da sie unabhängig von der Codierung sind, Bytes ist implementierungsdefiniert. Stringliterale haben den Typ string.

Operativer Betrieb

Bauchmuskeln

Semantik

Führt eine elementweise abs-Operation für den operand-Tensor durch und erzeugt einen result-Wert Tensor. Gehen Sie je nach Elementtyp so vor:

• Für vorzeichenbehaftete Ganzzahlen: Modulus ganzzahl.
• Für Gleitkommazahlen: abs aus IEEE-754.
• Für komplexe Zahlen: komplexer Modulus.
• Für quantisierte Typen: dequantize_op_quantize(abs, operand, type(result)).

Eingaben

Ausgaben

Einschränkungen

• (C1) shape(result) = shape(operand).
• (C2) baseline_element_type(result) ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • complex_element_type(element_type(operand)) wenn is_complex(operand).
  • Andernfalls baseline_element_type(operand).

Beispiele

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

Weitere Beispiele

Hinzufügen

Semantik

Führt eine elementweise Addition der beiden Tensoren lhs und rhs durch und erzeugt einen result-Tensor. Gehen Sie je nach Elementtyp so vor:

• Für boolesche Werte: logisches ODER.
• Für Ganzzahlen: Addition von Ganzzahlen.
• Für Gleitkommazahlen: addition aus IEEE-754.
• Für komplexe Zahlen: komplexe Addition
• Für quantisierte Typen: dequantize_op_quantize(add, lhs, rhs, type(result)).

Eingaben

Ausgaben

Einschränkungen

• Wenn die Operation nicht quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C1) type(lhs) = type(rhs) = type(result).
• Wenn die Operation quantisierte Tensoren verwendet: <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) Wenn is_per_axis_quantized(lhs), dann quantization_dimension(lhs) = quantization_dimension(result).
  • (C7) Wenn is_per_axis_quantized(rhs), dann quantization_dimension(rhs) = quantization_dimension(result).

Beispiele

// %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]]

Weitere Beispiele

after_all

Semantik

Stellt sicher, dass die Vorgänge, die die inputs erzeugen, vor jeder Ausführung ausgeführt werden Vorgänge, die von result abhängig sind. Bei der Ausführung dieses Vorgangs passiert nichts. Er existiert nur zum Einrichten von Datenabhängigkeiten von result bis inputs.

Eingaben

Ausgaben

Beispiele

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

Weitere Beispiele

all_gather

Semantik

Verkettet innerhalb jeder Prozessgruppe im StableHLO-Prozessraster die Werte der operands-Tensoren aus jedem Prozess entlang von all_gather_dim und erzeugt results-Tensoren.

Der Vorgang teilt das StableHLO-Prozessraster in process_groups auf, wie folgt definiert:

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

Danach gilt in jeder process_group Folgendes:

• operands...@receiver = [operand@sender for sender in process_group] für alle receiver in process_group.
• results...@process = concatenate(operands...@process, all_gather_dim) für alle process in process_group.

Eingaben

Ausgaben

Einschränkungen

• (C1) 0 <= all_gather_dim < rank(operands...).
• (C2) is_unique(replica_groups).
• (C3) size(replica_groups) ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas, wenn cross_replica verwendet wird.
  • num_replicas, wenn cross_replica_and_partition verwendet wird.
  • num_processes, wenn flattened_ids verwendet wird.
• (C4) 0 <= replica_groups < size(replica_groups).
• (C5) Wenn use_global_device_ids = true, dann channel_id > 0.
• (C6) type(results...) = type(operands...) mit Ausnahme von: <ph type="x-smartling-placeholder">
  </ph>
  • dim(results..., all_gather_dim) = dim(operands..., all_gather_dim) * dim(process_groups, 1).

Beispiele

// 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]]

Weitere Beispiele

all_reduce

Semantik

Wendet innerhalb jeder Prozessgruppe im StableHLO-Prozessraster eine Reduzierung an Funktion computation zu den Werten der operands-Tensoren aus jedem Prozess und erzeugt results-Tensoren.

Der Vorgang teilt das StableHLO-Prozessraster in process_groups auf, wie folgt definiert:

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

Danach gilt in jeder process_group Folgendes:

• results...@process[result_index] = exec(schedule) für einen binären Baum schedule, wobei gilt: <ph type="x-smartling-placeholder">
  </ph>
  • exec(node) = computation(exec(node.left), exec(node.right)).
  • exec(leaf) = leaf.value.
• schedule ist eine implementierungsdefinierte binäre Baumstruktur, deren Reihenfolge Durchlauf ist to_destination_type(operands...@process_group...[result_index], type(func_inputs(computation)[0])).

Eingaben

Ausgaben

Einschränkungen

• (C1) is_unique(replica_groups).
• (C2) size(replica_groups) ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas, wenn cross_replica verwendet wird.
  • num_replicas, wenn cross_replica_and_partition verwendet wird.
  • num_processes, wenn flattened_ids verwendet wird.
• (C3) 0 <= replica_groups < size(replica_groups).
• (C4) Wenn use_global_device_ids = true, dann channel_id > 0.
• (C5) computation hat den Typ (tensor<E>, tensor<E>) -> (tensor<E>), wobei is_promotable(element_type(operand), E)
• (C6) shape(results...) = shape(operands...).
• (C7) element_type(results...) = E.

Beispiele

// 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]

Weitere Beispiele

all_to_all

Semantik

Innerhalb jeder Prozessgruppe im StableHLO-Prozessraster werden die Werte der die operands-Tensoren entlang von split_dimension in Teile, verteilt die zwischen Prozessen und verkettet die verstreuten Teile entlang concat_dimension und erzeugt results-Tensoren. Der Vorgang teilt das StableHLO-Prozessraster in process_groups auf, wie folgt definiert:

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

Danach gilt in jeder process_group Folgendes:

• split_parts...@sender = split(operands...@sender, split_count, split_dimension) für alle sender in process_group.
• scattered_parts...@receiver = [split_parts...@sender[receiver_index] for sender in process_group], wobei receiver_index = process_group.index(receiver).
• results...@process = concatenate(scattered_parts...@process, concat_dimension).

Eingaben

Ausgaben

Einschränkungen

• (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) ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas, wenn cross_replica verwendet wird.
  • num_partitions, wenn cross_partition verwendet wird.
• (C7) 0 <= replica_groups < size(replica_groups).
• (C8) dim(replica_groups, 1) = split_count.
• (C9) type(results...) = type(operands...) mit Ausnahme von 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.

Beispiele

// 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]]

Weitere Beispiele

und

Semantik

Führt das elementweise UND der zwei Tensoren lhs und rhs durch und erzeugt eine result Tensor. Gehen Sie je nach Elementtyp so vor:

• Für boolesche Werte: logisches UND.
• Für Ganzzahlen: bitweises UND.

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]]

Weitere Beispiele

atan2

Semantik

Führt eine elementweise atan2-Operation für den Tensor lhs und rhs durch und erzeugt einen result-Tensor. Gehen Sie je nach Elementtyp so vor:

• Für Gleitkommazahlen: atan2 aus IEEE-754.
• Für komplexe Zahlen: komplexe atan2.
• Für quantisierte Typen: dequantize_op_quantize(atan2, lhs, rhs, type(result)).

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]

Weitere Beispiele

batch_norm_grad

Semantik

Berechnet Gradienten mehrerer Eingaben der Rückpropagierung von batch_norm_training aus grad_output und produziert grad_operand, grad_scale und grad_offset Tensoren. Formal kann diese Operation als eine Zerlegung in vorhandene StableHLO-Vorgänge mit der folgenden Python-Syntax:

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

Führt für quantisierte Typen durch 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))

Eingaben

Ausgaben

Einschränkungen

• (C1) 0 <= feature_index < rank(operand).
• (C2) operand, scale, mean, variance, grad_output, grad_operand, grad_scale und grad_offset haben dieselbe baseline_element_type.
• (C3) operand, grad_output und grad_operand haben die gleiche Form.
• (C4) scale, mean, variance, grad_scale und grad_offset haben die dieselbe Form haben.
• (C5) size(scale) = dim(operand, feature_index).

Beispiele

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

Semantik

Normalisiert den Tensor operand für alle Dimensionen mit Ausnahme der feature_index-Dimension und erzeugt einen result-Tensor. Formell gesprochen Operation kann als Zerlegung vorhandener StableHLO-Vorgänge ausgedrückt werden. mit der folgenden Python-Syntax:

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)

Führt für quantisierte Typen durch 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))

Eingaben

Ausgaben

Einschränkungen

• (C1) 0 <= feature_index < rank(operand).
• (C2) operand, scale, offset, mean, variance und result haben die Gleiche 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).

Beispiele

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

Semantik

Berechnet den Mittelwert und die Varianz für alle Dimensionen mit Ausnahme von feature_index Dimension und normalisiert den Tensor operand, der output, batch_mean erzeugt und batch_var-Tensoren. Formal kann dieser Vorgang als Zerlegung in vorhandene StableHLO-Vorgänge unter Verwendung der Python-Syntax als folgt:

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

Führt für quantisierte Typen durch 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))

Eingaben

Ausgaben

Einschränkungen

• (C1) 0 <= feature_index < rank(operand).
• (C2) Für operand, scale, offset, batch_mean, batch_var und output wurden denselben 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).

Beispiele

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

Semantik

Führt einen Bitcast-Vorgang für den operand-Tensor aus und erzeugt einen result-Tensor wobei die Bits des gesamten operand-Tensors mit der Methode Typ des result-Tensors.

Formeller gesprochen: E = element_type(operand), E' = element_type(result), und R = rank(operand):

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

bits gibt eine speicherinterne Darstellung eines bestimmten Werts und sein Verhalten zurück. ist implementierungsdefiniert, da die genaue Darstellung von Tensoren Implementierung definiert. Die genaue Darstellung von Elementtypen ist Implementierung definiert werden.

Eingaben

Ausgaben

Einschränkungen

• (C1) Für E = is_quantized(operand) ? storage_type(operand) : element_type(operand), E' = is_quantized(result) ? storage_type(result) : element_type(result) und R = rank(operand): <ph type="x-smartling-placeholder">
  </ph>
  • Wenn num_bits(E') = num_bits(E), shape(result) = shape(operand).
  • Wenn num_bits(E') < num_bits(E):
  • rank(result) = R + 1.
  • dim(result, i) = dim(operand, i) für alle 0 <= i < R.
  • dim(result, R) * num_bits(E') = num_bits(E).
  • Wenn num_bits(E') > num_bits(E):
  • rank(result) = R - 1.
  • dim(result, i) = dim(operand, i) für alle 0 <= i < R.
  • dim(operand, R - 1) * num_bits(E) = num_bits(E').
• (C2) Wenn is_complex(operand) or is_complex(result), dann is_complex(operand) and is_complex(result).

Beispiele

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

Weitere Beispiele

broadcast_in_dim

Semantik

Erweitert die Dimensionen und/oder den Rang eines Eingabetensors durch Duplizieren der Daten im operand-Tensor und erzeugt einen result-Tensor. Formeller gesprochen result[result_index] = operand[operand_index] – wo für alle d in axes(operand):

• operand_index[d] = 0 wenn dim(operand, d) = 1.
• Andernfalls operand_index[d] = result_index[broadcast_dimensions[d]].

Eingaben

Ausgaben

Einschränkungen

• (C1) element_type(result) ist gegeben durch: <ph type="x-smartling-placeholder">
  </ph>
  • element_type(operand), wenn !is_per_axis_quantized(operand).
  • element_type(operand) mit Ausnahme von quantization_dimension(operand), scales(operand) und zero_points(operand) können abweichen von quantization_dimension(result), scales(result) und zero_points(result) oder „Sonstiges“.
• (C2) size(broadcast_dimensions) = rank(operand).
• (C3) 0 <= broadcast_dimensions < rank(result).
• (C4) is_unique(broadcast_dimensions).
• (C5) Für alle d in axes(operand): <ph type="x-smartling-placeholder">
  </ph>
  • dim(operand, d) = 1 oder
  • dim(operand, d) = dim(result, broadcast_dimensions[d]).
• (C6) Wenn is_per_axis_quantized(result): <ph type="x-smartling-placeholder">
  </ph>
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)].
  • Wenn dim(operand, quantization_dimension(operand)) = 1, dann scales(result)[i] = scales(operand)[0] and zero_points(result)[i] = zero_points(operand)[0] for i in range(dim(result, quantization_dimension(result))).

Beispiele

// %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]
//            ]
//          ]

Weitere Beispiele

Supportanfrage

Semantik

Erzeugt die Ausgabe, wenn genau eine Funktion aus branches ausgeführt wird abhängig vom Wert von index. Formeller result = selected_branch() Dabei gilt:

• selected_branch = branches[index] wenn 0 <= index < size(branches).
• Andernfalls selected_branch = branches[-1].

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]

Weitere Beispiele

cbrt

Semantik

Führt eine elementweise kubische Wurzeloperation für den operand-Tensor durch und erzeugt eine result-Tensor. Gehen Sie je nach Elementtyp so vor:

• Für Gleitkommazahlen: rootn(x, 3) aus IEEE-754.
• Für komplexe Zahlen: komplexe kubische Wurzel.
• Für quantisierte Typen: dequantize_op_quantize(cbrt, operand, type(result))

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]

Weitere Beispiele

Ceil

Semantik

Führt die elementweise Obergrenze des operand-Tensors durch und erzeugt einen result-Tensor. Implementiert den roundToIntegralTowardPositive-Vorgang aus IEEE-754 Spezifikation zu ändern. Führt für quantisierte Typen durch dequantize_op_quantize(ceil, operand, type(result))

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]

Weitere Beispiele

Cholesky

Semantik

Berechnet die Cholesky-Zerlegung eines Batches von Matrizen.

Formeller gesprochen: für alle i in index_space(result), result[i0, ..., iR-3, :, :] ist eine Cholesky-Zerlegung von a[i0, ..., iR-3, :, :] in Form eines unteren Dreiecks (wenn lower true ist) oder eine obere dreieckige Form (wenn lower false ist) Matrix. Die Ausgabewerte im entgegengesetzten Dreieck, d.h. dem genauen oberen Dreieck oder sind implementierungsdefiniert.

Wenn es i gibt, bei dem die Eingabematrix keine hermitische positiv-definierte ist ist das Verhalten nicht definiert.

Führt für quantisierte Typen durch dequantize_op_quantize(lambda operand: cholesky(operand, lower), a, type(result))

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]
//          ]

einschränken

Semantik

Fixiert jedes Element des Tensors operand zwischen einem Mindest- und einem Höchstwert und erzeugt einen result-Tensor. Formeller, result[result_index] = minimum(maximum(operand[result_index], min_element), max_element), wobei min_element = rank(min) = 0 ? min[] : min[result_index], max_element = rank(max) = 0 ? max[] : max[result_index] Bei quantisierten Typen führt dequantize_op_quantize(clamp, min, operand, max, type(result)) aus.

Komplexe Zahlen müssen mit einer überraschenden Semantik, Daher planen wir, die Unterstützung für komplexe Zahlen in Zukunft einzustellen. für diesen Vorgang (#560).

Eingaben

Ausgaben

Einschränkungen

• (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).

Beispiele

// %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]

Weitere Beispiele

collective_broadcast

Semantik

Innerhalb jeder Prozessgruppe im StableHLO-Prozessraster senden Sie den Wert des operand-Tensor vom Quellprozess zu den Zielprozessen und erzeugt einen result-Tensor.

Der Vorgang teilt das StableHLO-Prozessraster in process_groups auf, wie folgt definiert:

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

Danach wird result@process folgendermaßen angegeben:

• operand@process_groups[i, 0], wenn ein i vorhanden ist, sodass der Prozess in process_groups[i].
• broadcast_in_dim(constant(is_quantized(result) ? quantize(0, element_type(result)) : 0, element_type(result)), [], type(result)) sonst.

Eingaben

Ausgaben

Einschränkungen

• (C1) is_unique(replica_groups).
• (C2) 0 <= replica_groups < N, wobei N definiert ist als: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas, wenn cross_replica verwendet wird.
  • num_partitions, wenn cross_partition verwendet wird.
• (C3) type(result) = type(operand).

Beispiele

// 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

Semantik

Innerhalb jeder Prozessgruppe im StableHLO-Prozessraster wird der Wert der operand-Tensor vom Quellprozess zum Zielprozess und erzeugt einen result-Tensor.

Der Vorgang teilt das StableHLO-Prozessraster in process_groups auf, wie folgt definiert:

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

Danach wird result@process folgendermaßen angegeben:

• operand@process_groups[i, 0], wenn ein i vorhanden ist, sodass process_groups[i, 1] = process
• broadcast_in_dim(constant(is_quantized(result) ? quantize(0, element_type(result)) : 0, element_type(result)), [], type(result)) sonst.

Eingaben

Ausgaben

Einschränkungen

• (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, wobei N definiert ist als: <ph type="x-smartling-placeholder">
  </ph>
  • num_replicas, wenn cross_replica verwendet wird.
  • num_partitions, wenn cross_partition verwendet wird.
• (C5) type(result) = type(operand).

Beispiele

// 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]]

Weitere Beispiele

vergleichen

Semantik

Führt einen elementweisen Vergleich von lhs- und rhs-Tensoren durch comparison_direction und compare_type und erzeugt einen result-Tensor.

Für die Werte comparison_direction und compare_type gilt Folgendes: Semantik:

Für boolesche und ganzzahlige Elementtypen:

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

Für Gleitkommaelementtypen mit compare_type = FLOAT wird durch die Operation implementiert folgende IEEE-754-Vorgänge:

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

Bei Gleitkomma-Elementtypen mit compare_type = TOTALORDER wird die Operation verwendet die Kombination aus totalOrder- und compareQuietEqual-Vorgängen aus IEEE-754.

Bei komplexen Elementtypen ist der lexikografische Vergleich von (real, imag)-Paaren wird mit den angegebenen comparison_direction und compare_type ausgeführt. Komplexe Zahlen müssen mit einer überraschenden Semantik, Daher planen wir, die Unterstützung für komplexe Zahlen in Zukunft einzustellen. wenn comparison_direction GE, GT, LE oder LT ist (#560)

Für quantisierte Typen. führt dequantize_compare(lhs, rhs, comparison_direction) aus.

Eingaben

Ausgaben

Einschränkungen

• (C1) baseline_element_type(lhs) = baseline_element_type(rhs).
• (C2) shape(lhs) = shape(rhs) = shape(result).
• (C3) compare_type ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • SIGNED wenn is_signed_integer(element_type(lhs)).
  • UNSIGNED wenn is_unsigned_integer(element_type(lhs)) or is_boolean(element_type(lhs)).
  • FLOAT oder TOTALORDER, wenn is_float(element_type(lhs)).
  • FLOAT wenn is_complex(element_type(lhs)).

Beispiele

// %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]

Weitere Beispiele

komplex

Semantik

Führt eine elementweise Konvertierung in einen komplexen Wert aus einem Paar aus reellen und lhs und rhs und erzeugt einen result-Tensor.

Eingaben

Ausgaben

Einschränkungen

• (C1) type(lhs) = type(rhs).
• (C2) shape(result) = shape(lhs).
• (C3) element_type(result) hat den Typ complex<E>, wobei E = element_type(lhs)

Beispiele

// %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)]

Weitere Beispiele

Zusammengesetzt

Semantik

Kapselt einen Vorgang ein, der aus anderen StableHLO-Vorgängen besteht, inputs und composite_attributes nehmen und dabei results produzieren. Die Die Semantik des Vorgangs wird durch das Attribut decomposition implementiert. Die composite-Op kann durch seine Zerlegung ersetzt werden, ohne das Programm zu ändern Semantik. In Fällen, in denen die Inline-Zerlegung nicht die gleichen Op-Semantik. Verwenden Sie am besten custom_call.

Mit dem Feld version (standardmäßig 0) wird angegeben, wann die eine Änderung der Semantik.

Eingaben

Ausgaben

Einschränkungen

• (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)

Beispiele

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

Weitere Beispiele

concatenate

Semantik

Verkettet inputs entlang der Dimension dimension in derselben Reihenfolge wie bei und erzeugt einen result-Tensor. Formeller gesprochen result[i0, ..., id, ..., iR-1] = inputs[k][i0, ..., kd, ..., iR-1]. Dabei gilt:

1. id = d0 + ... + dk-1 + kd.
2. d ist gleich dimension und d0 sind die d. Dimensionsgrößen von inputs.

Eingaben

Ausgaben

Einschränkungen

• (C1) same(element_type(inputs...)).
• (C2) same(shape(inputs...)) mit Ausnahme von 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]) mit Ausnahme von: <ph type="x-smartling-placeholder">
  </ph>
  • dim(result, dimension) = dim(inputs[0], dimension) + ....

Beispiele

// %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]]

Weitere Beispiele

Konstante

Semantik

Erzeugt einen output-Tensor aus einer konstanten value.

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

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

Weitere Beispiele

eine Conversion ausführen

Semantik

Führt eine elementweise Konvertierung von einem Elementtyp in einen anderen durch operand-Tensor und erzeugt einen result-Tensor.

Für Conversions vom Typ boolean-to-any-supported-type ist der Wert false in Null und der Wert true in eins konvertiert. Für any-supported-type-to-boolean-Conversions wird der Wert Null in false und Werte ungleich null werden in true konvertiert. Unten sehen Sie, wie dies auch für komplexe Typen.

Für Konvertierungen mit integer-to-integer und integer-to-floating-point oder floating-point-to-floating-point, wenn der Quellwert genau im Zieltyp dargestellt, ist der Ergebniswert genau Darstellung. Ansonsten steht das Verhalten noch nicht fest. (#180)

Bei Konvertierungen mit floating-point-to-integer ist der Bruchteil abgeschnitten. Wenn der abgeschnittene Wert nicht im Zieltyp dargestellt werden kann, Das Verhalten steht noch nicht fest (#180).

Konvertierungen, die komplex zu komplex umfassen, folgen dem gleichen Verhalten wie floating-point-to-floating-point-Conversions, mit denen reale und imaginären Teilen.

Für Conversions vom Typ complex-to-any-other-type und complex-to-any-other-type der imaginäre Quellwert wird ignoriert oder der imaginäre Zielwert ist auf 0 gesetzt. Die Umrechnung des realen Teils folgt der Gleitkomma-Conversions.

Prinzipiell könnte diese Operation eine Dequantisierung (Konvertierung von quantisierte Tensoren in reguläre Tensoren), Quantisierung (Umwandlung von regulären Tensoren) Tensoren in quantisierte Tensoren) und Requantisierung (Umwandlung zwischen quantisierten Tensoren Tensoren), aber derzeit haben wir dafür dedizierte Operationen – uniform_dequantize für den ersten Anwendungsfall und uniform_quantize für den zweiten und dritten Anwendungsfälle. In Zukunft werden diese beiden Vorgänge möglicherweise zusammengeführt. in convert (#1576) ein.

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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)]

Weitere Beispiele

Faltung

Semantik

Berechnet Punktprodukte zwischen Fenstern von lhs und Segmenten von rhs und erzeugt result Das folgende Diagramm zeigt, wie Elemente in result aus lhs und rhs anhand eines konkreten Beispiels.

Formaler sollten Sie die folgende Neuformulierung der Eingaben in Bezug auf lhs betrachten: um Fenster von lhs ausdrücken zu können:

• 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).

Für diese Neuausrichtung werden die folgenden Hilfsfunktionen verwendet:

• 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], wobei j[d] = i[permutation[d]].

Wenn feature_group_count = 1 und batch_group_count = 1, dann für alle output_spatial_index in index_space(dim(result, output_spatial_dimensions...)), result[result_shape(:, output_spatial_index, :)] = dot_product, wobei:

• 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]) Diese Funktion wird offenbar noch nicht genutzt. Daher planen wir, sie in Zukunft zu entfernen. es (#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]).

Wenn 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).

Wenn 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)

Führt für quantisierte Typen 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)) aus.

Führt für hybrid quantisierte Typen 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) aus.

Eingaben

Ausgaben

Einschränkungen

• (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) Angegebener 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) Angegebener 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) Angegebener 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) ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • dim(lhs, input_batch_dimension) / batch_group_count wenn result_dim = output_batch_dimension.
  • dim(rhs, kernel_output_feature_dimension) wenn result_dim = output_feature_dimension.
  • num_windows, andernfalls gilt:
  • 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.
• Wenn die Operation nicht quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C27) element_type(lhs) = element_type(rhs) = element_type(result).
• Wenn die Operation quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C28) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs).
  • (C29) Wenn is_per_axis_quantized(rhs), danach quantization_dimension(rhs) = kernel_output_feature_dimension.
  • (C30) Wenn is_per_axis_quantized(result), dann quantization_dimension(result) = output_feature_dimension.
  • Wenn is_quantized(lhs):
  • (C31) storage_type(lhs) = storage_type(rhs).
  • (C32) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C33) Wenn is_per_tensor_quantized(rhs), dann is_per_tensor_quantized(result).
  • Wenn !is_quantized(lhs):
  • (C34) element_type(lhs) = expressed_type(rhs) = element_type(result).

Beispiele

// %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]]
//          ]]

Weitere Beispiele

Kosinus

Semantik

Führt eine elementweise Kosinusoperation für den operand-Tensor durch und erzeugt einen result-Tensor. Gehen Sie je nach Elementtyp so vor:

• Für Gleitkommazahlen: cos aus IEEE-754.
• Für komplexe Zahlen: komplexer Kosinus.
• Für quantisierte Typen: dequantize_op_quantize(cosine, operand, type(result)).

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]]

Weitere Beispiele

count_leading_zeros

Semantik

Führt eine elementweise Zählung der Anzahl der führenden Null-Bits im operand durch und erzeugt einen result-Tensor.

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

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

Weitere Beispiele

custom_call

Semantik

Kapselt einen implementierungsdefinierten Vorgang call_target_name, der inputs und called_computations und produziert results. has_side_effect, backend_config und api_version können verwendet werden, um zusätzliche implementierungsdefinierten Metadaten.

Im Moment enthält dieser Vorgang eine ziemlich unorganisierte Sammlung von die die organische Entwicklung des entsprechenden Betriebs in den XLA-Compiler. Wir planen, diese Metadaten in Zukunft zu vereinheitlichen. (#741)

Eingaben

Ausgaben

Beispiele

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

Dividieren

Semantik

Führt die elementweise Division der Tensoren des Dividenden lhs und des Divisors rhs durch und erzeugt einen result-Tensor. Gehen Sie je nach Elementtyp so vor:

• Für Ganzzahlen: Ganzzahldivision, die den algebraischen Quotienten mit einer beliebigen Bruchteil verworfen.
• Für Gleitkommazahlen: division aus IEEE-754.
• Für komplexe Zahlen: komplexe Division.
• Für quantisierte Typen: <ph type="x-smartling-placeholder">
  </ph>
  • dequantize_op_quantize(divide, lhs, rhs, type(result)).

Eingaben

Ausgaben

Einschränkungen

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

Beispiele

// %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]

Weitere Beispiele

dot_general

Semantik

Berechnet die Punktprodukte zwischen Segmenten von lhs und Segmenten von rhs und erzeugt einen result-Tensor.

Formeller result[result_index] = dot_product, wobei Folgendes gilt:

• 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 wobei size(result_batching_index) = size(lhs_batching_dimensions), size(result_lhs_index) = size(lhs_result_dimensions) und 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))

Führt für quantisierte Typen 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)) aus.

Führt für hybrid quantisierte Typen 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) aus.

precision_config steuert den Kompromiss zwischen Geschwindigkeit und Genauigkeit auf Beschleuniger-Back-Ends. Dies kann einer der folgenden sein (im ist die Semantik dieser enum-Werte zu niedrig angegeben. wie Sie dieses Problem #755):

• DEFAULT: Schnellste Berechnung, aber am wenigsten genaue Schätzung zum die ursprüngliche Nummer.
• HIGH: langsamere Berechnung, aber genauere Annäherung an den Wert die ursprüngliche Nummer.
• HIGHEST: Die langsamste Berechnung, aber die genaueste Annäherung an den Wert die ursprüngliche Nummer.

Ein DotAlgorithm definiert die Haupteigenschaften des Algorithmus, der zur Implementierung verwendet wird. die Punktoperation, die auch die Genauigkeit definiert. Wenn das Algorithmusattribut festgelegt sind, muss precision_config auf DEFAULT gesetzt sein. DotAlgorithms haben keinen Standardwert, da die Standardparameter definiert. Daher können alle Punktalgorithmusfelder auf None gesetzt werden, um ein Algorithmus mit leerem Punkt, der stattdessen den Wert precision_config verwendet.

DotAlgorithm-Felder enthalten:

• lhs_precision_type und rhs_precision_type, die Genauigkeit, die die linke und die linke Seite und die Die RHS-Werte des Vorgangs werden auf gerundet. Genauigkeitstypen sind unabhängig vom die Speichertypen der Ein- und Ausgabe.
• accumulation_type ist die für die Akkumulation verwendete Genauigkeit.
• lhs_component_count, rhs_component_count und num_primitive_operations anwenden, wenn wir einen Algorithmus ausführen, der die linke und/oder rechte Seite in und führt mehrere "primitive" Punkt-Operationen auf diesen -Werte haben, um eine höhere Genauigkeit zu emulieren (z.B. Mit dem Datentyp bfloat16 „Artificial Intelligence“ für präzisere Berechnungen arbeiten: bf16_6x tf32_3x usw.). Bei Algorithmen ohne Zerlegung sind diese Werte sollte auf 1 festgelegt sein.
• allow_imprecise_accumulation, um anzugeben, ob die Akkumulation mit geringerer Genauigkeit verwendet wird ist für einige Schritte zulässig (z.B. CUBLASLT_MATMUL_DESC_FAST_ACCUM).

Beispiele für DotAlgorithm-Attribute:

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

Es liegt an den Implementierungen, zu entscheiden, welche Kombinationen unterstützt werden. In kann nicht garantiert werden, dass jeder Algorithmus Beschleunigertyp durch den Nutzer des StableHLO. Wenn ein bestimmter Algorithmus sollte ein Fehler ausgegeben werden, anstatt auf einen als Alternative. Die StableHLO-Überprüfung ermöglicht eine bestmögliche Verifizierung, um Algorithmen zu verhindern, die bekanntermaßen auf keiner Hardware unterstützt werden.

Weitere Informationen finden Sie unter xla_data.proto > Algorithm. für einige unterstützte Algorithmuswerte. Ticket Nr. 2483 enthält den Plan zur Erstellung eines zentralisiertes Dokument zu unterstützten Algorithmen vom Back-End.

Eingaben

Ausgaben

Einschränkungen

• (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).
• Wenn die Operation nicht quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C13) element_type(lhs) = element_type(rhs).
• Wenn die Operation quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C14) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs).
  • (C15) zero_points(rhs) = 0.
  • (C16) Wenn is_per_axis_quantized(rhs), dann quantization_dimension(rhs) nicht in rhs_contracting_dimensions.
  • Wenn is_quantized(lhs):
  • (C17) storage_type(lhs) = storage_type(rhs).
  • (C18) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C19) Wenn is_per_tensor_quantized(rhs), dann is_per_tensor_quantized(result).
  • Wenn !is_quantized(lhs):
  • (C20) element_type(lhs) = expressed_type(rhs) = element_type(result).
• Wenn !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.

Beispiele

// %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]]
//          ]

Weitere Beispiele

dynamic_broadcast_in_dim

Semantik

Funktional ist dieser Vorgang identisch mit broadcast_in_dim op, aber die Ergebnisform wird dynamisch über output_dimensions angegeben.

Für den Vorgang werden auch die optionalen Attribute known_expanding_dimensions und known_non_expanding_dimensions akzeptiert. um statisches Wissen über das Expansionsverhalten von Dimensionen auszudrücken. Wenn nicht angegeben, wird angenommen, dass alle Abmessungen möglicherweise maximiert werden.

Eingaben

Ausgaben

Einschränkungen

• (C1) element_type(result) ist gegeben durch: <ph type="x-smartling-placeholder">
  </ph>
  • element_type(operand), wenn !is_per_axis_quantized(operand).
  • element_type(operand) mit Ausnahme von quantization_dimension(operand), scales(operand) und zero_points(operand) können abweichen von quantization_dimension(result), scales(result) und zero_points(result) oder „Sonstiges“.
• (C2) size(broadcast_dimensions) = rank(operand).
• (C3) 0 <= broadcast_dimensions < rank(result).
• (C4) is_unique(broadcast_dimensions).
• (C5) Für alle d in axes(operand): <ph type="x-smartling-placeholder">
  </ph>
  • dim(operand, d) = 1 oder
  • dim(operand, d) = dim(result, broadcast_dimensions[d]).
• (C6) Wenn is_per_axis_quantized(result): <ph type="x-smartling-placeholder">
  </ph>
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)].
  • Wenn dim(operand, quantization_dimension(operand)) = 1, dann 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).

Beispiele

// %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]
//            ]
//          ]

Weitere Beispiele

dynamic_conv

Semantik

Funktional ist dieser Vorgang identisch mit Faltung op, aber das Padding wird dynamisch über padding angegeben.

Eingaben

Ausgaben

Einschränkungen

• (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) Angegebener 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) Angegebener 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) Angegebener 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) ist definiert als: <ph type="x-smartling-placeholder">
  </ph>
  • dim(lhs, input_batch_dimension) / batch_group_count wenn result_dim = output_batch_dimension.
  • dim(rhs, kernel_output_feature_dimension) wenn result_dim = output_feature_dimension.
  • num_windows, andernfalls gilt:
  • 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.
• Wenn die Operation nicht quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C27) element_type(lhs) = element_type(rhs) = element_type(result).
• Wenn die Operation quantisierte Tensoren verwendet: <ph type="x-smartling-placeholder">
  </ph>
  • (C28) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs).
  • (C29) Wenn is_per_axis_quantized(rhs), danach quantization_dimension(rhs) = kernel_output_feature_dimension.
  • (C30) Wenn is_per_axis_quantized(result), dann quantization_dimension(result) = output_feature_dimension.
  • Wenn is_quantized(lhs):
  • (C31) storage_type(lhs) = storage_type(rhs).
  • (C32) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result).
  • (C33) Wenn is_per_tensor_quantized(rhs), dann is_per_tensor_quantized(result).
  • Wenn !is_quantized(lhs):
  • (C34) element_type(lhs) = expressed_type(rhs) = element_type(result).

Beispiele

// %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]]
//          ]]