StableHLO 規格

StableHLO 是機器學習 (ML) 模型中高階運算 (HLO) 的運算集。StableHLO 可做為不同 ML 架構和 ML 編譯器之間的可攜性層：產生 StableHLO 程式的 ML 架構與使用 StableHLO 程式的 ML 編譯器相容。

我們的目標是透過在各種機器學習架構 (例如 TensorFlow、JAX 和 PyTorch) 與機器學習編譯器 (例如 XLA 和 IREE) 之間建立更多互通性，簡化機器學習開發流程並加快開發速度。為此，本文件提供 StableHLO 程式設計語言的規範。

這個規格包含三個主要部分首先，程式一節說明 StableHLO 程式的結構，其中包含 StableHLO 函式，而這些函式本身則包含 StableHLO 運算。在該結構中，「Ops」部分會指定個別操作的語意。「Execution」部分會為在程式中一起執行的所有運算提供語意。最後，符號一節會討論規格中使用的符號。

如要查看 StableHLO 先前版本的規格，請開啟您感興趣的標記版本的存放區。例如 StableHLO v0.19.0 規格。如要查看 StableHLO 在每個子版本遞增發生的變更，請參閱 VhloDialect.td 中的版本記錄。

程式

Program ::= {Func}

StableHLO 程式包含任意數量的 StableHLO 函式。以下是含有函式 @main 的程式範例，該函式有 3 個輸入 (%image、%weights 和 %bias) 和 1 個輸出。函式主體有 6 個運算。

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

函式

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

StableHLO 函式 (也稱為已命名函式) 包含一個 ID、輸入/輸出項目和主體。我們計劃日後為函式推出其他中繼資料，以提升與 HLO 的相容性 (#425、#626、#740、#744)。

ID

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

StableHLO ID 與許多程式語言中的 ID 類似，但有兩個特點：1) 所有 ID 都有符號，用於區分不同類型的 ID；2) 值 ID 可以完全以數字表示，簡化 StableHLO 程式的產生作業。

類型

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

StableHLO 類型分為值類型 (也稱為第一類型)，代表 StableHLO 值，以及非值類型，用於描述其他程式元素。StableHLO 類型與許多程式設計語言中的類型相似，主要特徵是 StableHLO 的特定領域性質，會產生一些不尋常的結果 (例如，純量類型不是值類型)。

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

張量類型代表張量，也就是多維陣列。包含形狀和元素類型，其中形狀代表依據對應維度 (也稱為「ax」) 的遞增順序，從 0 到 R-1 的遞增順序表示非負數或不明維度大小。維度 R 的數量稱為「排名」。舉例來說，tensor<2x3xf32> 是形狀為 2x3 的張量類型，元素類型為 f32。它有兩個維度 (或稱兩個軸) - 第 0 維度和第 1 維度，大小分別為 2 和 3。其排名為 2。

形狀可能部分或完全不明 (動態)，例如 tensor<?x2xf64> 部分不明，而 tensor<?x?xf64> 完全不明。動態維度大小會使用 ? 表示。形狀無法排序。

我們未來會嘗試將張量類型擴展至維度大小和元素類型以外的其他類型，例如納入版面配置 (#629) 和稀疏度 (#1078)。

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

量化元素類型代表儲存類型的整數值，範圍為 storage_min 到 storage_max (含)，對應於表示型別的浮點值。對於指定的整數值 i，對應的浮點值 f 可計算為 f = (i - zero_point) * scale，其中 scale 和 zero_point 稱為量化參數。storage_min 和 storage_max 是文法中的選用項目，但預設值為 min_value(storage_type) 和 max_value(storage_type)。量化元素類型具有下列限制：

• (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) 如果 is_empty(quantization_dimension)，則 size(scales) = 1。
• (C11) 0 <= quantization_dimension。

目前 QuantizationScale 是浮點常數，但使用者對以整數為基礎的比例表達方式 (以乘數和位移表示) 非常感興趣。我們預計在近期內探索這項功能。(#1404)。

我們目前正在討論 QuantizationZeroPoint 的語意，包括類型、值，以及在量化張量類型中是否只能有一個或多個零點。根據這次討論的結果，零點的規格日後可能會改變 (#1405)。

另一個正在討論的議題是 QuantizationStorageMin 和 QuantizationStorageMax 的語意，以決定是否應對這些值和量化張量值施加任何限制 (#1406)。

最後，我們計劃探索如何呈現未知的比例和零點，就像我們計劃探索如何呈現未知的維度大小一樣 (#1407)。

量化張量類型代表具有量化元素的張量。這些張量與一般張量完全相同，但元素具有量化元素類型，而非一般元素類型。

在量化張量中，量化可以是「每個張量」，也就是為整個張量提供一個 scale 和 zero_point；也可以是「每個軸」，也就是為特定維度 quantization_dimension 的每個切片提供一組 scales 和 zero_points。更正式來說，在張量 t 中，每軸量化會有 quantization_dimension 的 dim(t, quantization_dimension) 配量：t[:, ..., 0, ..., :], t[:, ..., 1, ..., :] 等。i 切片中的所有元素都使用 scales[i] 和 zero_points[i] 做為量化參數。量化張量類型有下列限制：

• 針對每個張量量化：
  • 沒有其他限制。
• 針對每個軸的量化：
  • (C12) quantization_dimension < rank(self)。
  • (C13) dim(self, quantization_dimension) = size(scales)。

TokenType ::= 'token'

符記類型代表符記，也就是某些作業產生及消耗的非公開值。符記可用來為作業施加執行順序，如執行一節所述。

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

元組類型代表元組，也就是異質清單。元組是舊版功能，僅存在於 HLO 相容性。在 HLO 中，元組用於表示可變參數輸入和輸出。在 StableHLO 中，變數輸入和輸出是原生支援的，而 StableHLO 中唯一的用途是全面代表 HLO ABI，例如 T、tuple<T> 和 tuple<tuple<T>> 可能會因特定實作而有實質差異。日後，我們預計會對 HLO ABI 進行變更，這可能會讓我們從 StableHLO 中移除元組類型 (#598)。

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

元素類型代表張量類型的元素。與許多程式設計語言不同，這些類型並不是 StableHLO 中的第一類別。這表示 StableHLO 程式無法直接表示這些類型的值 (因此，使用 tensor<T> 類型的 0 維度張量值，是慣用 T 類型的純量值)。

• 布林值類型代表布林值 true 和 false。
• 整數類型可以是帶正負號 (si) 或不帶正負號 (ui)，並且具有支援的位元寬度 (2、4、8、16、32 或 64)。帶正負號的 siN 類型代表從 -2^(N-1) 到 2^(N-1)-1 的整數值 (含括在內)，而未帶正負號的 uiN 類型則代表從 0 到 2^N-1 的整數值 (含括在內)。
• 浮點類型可以是下列任一值：
  • f8E3M4、f8E4M3 和 f8E5M2 8 位元浮點數，遵循 IEEE-754 慣例。
  • f8E4M3FN 和 f8E5M2 類型分別對應至 深度學習的 FP8 格式中所述 FP8 格式的 E4M3 和 E5M2 編碼。
  • 與 E4M3 和 E5M2 格式相對應的 f8E4M3FNUZ 和 f8E5M2FNUZ 類型，如深層類神經網路的 8 位元數值格式中所述。
  • f8E4M3B11FNUZ 類型對應至 E4M3 編碼，如 Hybrid 8-bit Floating Point (HFP8) Training and Inference for Deep Neural Networks 所述。
  • bf16 類型，對應至 BFloat16：在 Cloud TPU 上取得高效能的秘訣一文中所述的 bfloat16 格式。
  • f16、f32 和 f64 類型分別對應於 binary16 (「半精度」)、binary32 (「單精度」) 和 binary64 (「雙精度」) 格式，如 IEEE 754 標準中所述。
  • tf32 類型對應的是 TensorFloat32 格式，且對 StableHLO 的支援有限。
  • OCP Microscaling 格式規格中所述的 f4E2M1FN、f6E2M3FN、f6E3M2FN 和 f8E8M0FNU MX (微縮放) 類型。
• 複雜類型代表包含相同元素類型「實際部分」和「虛部分」的複雜值。支援的複合類型為 complex<f32> (兩個部分皆為 f32 類型) 和 complex<f64> (兩個部分皆為 f64 類型)。

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

函式類型代表已命名和匿名函式。這些資料類型具有輸入類型 (-> 左側的類型清單) 和輸出類型 (-> 右側的類型清單)。在許多程式設計語言中，函式類型是第一類，但在 StableHLO 中則不是。

StringType ::= 'string'

字串型別代表位元組序列。與許多程式語言不同，字串類型在 StableHLO 中並非第一類型，只用於為程式元素指定靜態中繼資料。

作業

StableHLO 運算 (也稱為 ops) 代表機器學習模型中一組封閉的高階運算。如上所述，StableHLO 語法受到 MLIR 的啟發，這不一定是最符合人體工學的替代方案，但最適合 StableHLO 的目標：在機器學習架構和機器學習編譯器之間建立更高的互通性。

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

StableHLO 作業 (也稱為 ops) 具有名稱、輸入/輸出和簽章。名稱包含 stablehlo. 前置字元和助憶法，可用於唯一識別其中一個支援的作業。請參閱下列完整清單，瞭解所有支援的作業。

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

作業會消耗輸入並產生輸出。輸入內容分為輸入值 (在執行期間計算)、輸入函式 (以靜態方式提供，因為在 StableHLO 中，函式不是一級值) 和輸入屬性 (同樣以靜態方式提供)。運算消耗及產生的輸入和輸出種類取決於其記憶法。舉例來說，add 運算子會使用 2 個輸入值，並產生 1 個輸出值。相比之下，select_and_scatter 運算會使用 3 個輸入值、2 個輸入函式和 3 個輸入屬性。

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

輸入函式 (也稱為「匿名函式」) 與已命名函式非常相似，但有以下差異：1) 沒有 ID (因此稱為「匿名」)，2) 不會宣告輸出類型 (輸出類型會從函式中的 return op 推斷而來)。

輸入函式的語法包含目前未使用的部分 (請參閱上述的 Unused 正式環境)，這是能與 MLIR 相容性的部分。在 MLIR 中，有一個更廣義的「區域」概念，可透過跳躍運算子連結多個「運算子」的「區塊」。這些區塊的 ID 會對應至 Unused 產出作業，因此可彼此區分。StableHLO 沒有跳躍運算，因此 MLIR 語法的對應部分未使用 (但仍存在)。

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

輸入屬性具有名稱和值，其中值為支援的常數之一。這是為節目元素指定靜態中繼資料的主要方式。舉例來說，concatenate 運算子會使用屬性 dimension 指定輸入值連接的維度。同樣地，slice 運算子會使用 start_indices 和 limit_indices 等多個屬性，指定用於切片輸入值的邊界。

目前，實際運作的 StableHLO 程式有時會包含本文件未提及的屬性。我們計劃日後將這些屬性吸收至 StableHLO 運算元，或禁止這些屬性出現在 StableHLO 程式中。在此期間，請參閱下列屬性清單：

• layout (#629)。
• mhlo.frontend_attributes (#628)。
• mhlo.sharding (#619)。
• output_operand_aliases (#740)。
• 地點中繼資料 (#594)。

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

運算子簽章包含所有輸入值的類型 (-> 左側的類型清單) 和所有輸出值的類型 (-> 右側的類型清單)。嚴格來說，輸入類型是多餘的，輸出類型也幾乎總是多餘的 (因為對於大多數 StableHLO 運算子，輸出類型可從輸入推斷)。不過，為了與 MLIR 相容，我們刻意將運算子簽章納入 StableHLO 語法。

以下是其記憶法為 select_and_scatter 的運算範例。這個函式會使用 3 個輸入值 (%operand、%source 和 %init_value)、2 個輸入函式和 3 個輸入屬性 (window_dimensions、window_strides 和 padding)。請注意，這個函式的簽章只包含輸入值的類型 (但不包含內嵌提供的輸入函式和屬性類型)。

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

常數

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

StableHLO 常數具有常值和類型，兩者共同代表 StableHLO 值。一般來說，類型是常數語法的一部分，除非是不混淆的情況 (例如布林值常數具有 i1 類型，而整數常數可以有多個可能的類型)。

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

布林值常數代表布林值 true 和 false。布林常數具有 i1 類型。

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

整數常數會透過使用十六進制標記法的字串表示整數值。系統不支援其他基數，例如二進位或八進位。整數常數有下列限制：

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

浮點常數會透過使用十進位或科學記號的字串表示浮點值。此外，十六進位記號法可用於直接指定對應類型的浮點格式中的基本位元。浮點常數具有下列限制：

• (C1) 如果使用非十六進位標記法，is_wellformed(float_literal, float_type)。
• (C2) 如果使用十六進位表示法，則為 size(hexadecimal_digits) = num_bits(float_type) / 4。

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

複數常數會使用實數部分 (先出現) 和虛數部分 (後出現) 的清單，代表複數值。例如，(1.0, 0.0) : complex<f32> 代表 1.0 + 0.0i，(0.0, 1.0) : complex<f32> 代表 0.0 + 1.0i。這些部分儲存在記憶體中的順序是由實作定義。複雜常數有下列限制：

• (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 常數使用 NumPy 標記法指定的巢狀清單來代表張量值。舉例來說，dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> 代表張量值，將以下從索引對應至元素：{0, 0} => 1、{0, 1} => 2、{0, 2} => 3、{1, 0} => 4、{1, 1} => 5、{1, 2} => 6。這些元素在記憶體中儲存的順序由實作方式定義。張量常數有下列限制：

• (C1) has_syntax(tensor_literal, element_type(tensor_type))，其中：
  • 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))，其中：
  • 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:])。
  • 否則為 false。

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

量化張量常數會使用與張量常數相同的符號表示量化張量值，其中元素會指定為其儲存類型的常數。量化張量常數有下列限制：

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

字串常值由使用 ASCII 字元和逸出序列指定的位元組所組成。這些值不受編碼方式影響，因此這些位元組的解讀方式由實作方式定義。字串常值的類型為 string。

作業數

腹部

語義學

對 operand 張量執行元素級別的絕對值運算，並產生 result 張量。視元素類型而定，執行下列操作：

• 如為帶號整數：整數模數。
• 浮點值：IEEE-754 的 abs。
• 針對複數：複數的模數。
• 對於量化類型：dequantize_op_quantize(abs, operand, type(result))。

輸入

輸出內容

限制

• (C1) shape(result) = shape(operand)。
• (C2) baseline_element_type(result) 的定義為：
  • 若 is_complex(operand) 則設為 complex_element_type(element_type(operand))。
  • 否則為 baseline_element_type(operand)。

範例

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

更多範例

add

語義學

這個外掛程式能執行兩個張量 lhs 和 rhs 的元素相關新增，並產生 result 張量。視元素類型而定，執行下列操作：

• 布林值：邏輯 OR。
• 整數：加整數。
• 浮點值：IEEE-754 的 addition。
• 複數：複數加法。
• 對於量化類型：dequantize_op_quantize(add, lhs, rhs, type(result))。

輸入

輸出內容

限制

• 如果運算使用非量化的張量：
  • (C1) type(lhs) = type(rhs) = type(result)。
• 如果運算使用經過量化的張量：
  • (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) 如果 is_per_axis_quantized(lhs)，則 quantization_dimension(lhs) = quantization_dimension(result)。
  • (C7) 如果 is_per_axis_quantized(rhs)，則 quantization_dimension(rhs) = quantization_dimension(result)。

範例

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

更多範例

after_all

語義學

確保產生 inputs 的作業會在任何依賴 result 的作業之前執行。這項作業不執行任何動作，只是為了建立從 result 到 inputs 的資料依附元件。

輸入

輸出內容

範例

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

更多範例

all_gather

語義學

在 StableHLO 程序格線的每個程序群組中，沿著 all_gather_dim 連結每個程序的 operands 張量值，並產生 results 張量。

這項作業會將 StableHLO 處理格線分割為 process_groups，其定義如下：

• cross_replica(replica_groups) 若 channel_id <= 0 and use_global_device_ids = false。
• 如果為 channel_id > 0 and use_global_device_ids = false，則為 cross_replica_and_partition(replica_groups)。
• flattened_ids(replica_groups) 若 channel_id > 0 and use_global_device_ids = true。

接著，在每個 process_group 中：

• process_group 中的所有 receiver 的 operands...@receiver = [operand@sender for sender in process_group]。
• process_group 中的所有 process 的 results...@process = concatenate(operands...@process, all_gather_dim)。

輸入

輸出內容

限制

• (C1) 0 <= all_gather_dim < rank(operands...)。
• (C2) is_unique(replica_groups)。
• (C3) size(replica_groups) 的定義為：
  • 如果使用 cross_replica，則為 num_replicas。
  • 如果使用 cross_replica_and_partition，則為 num_replicas。
  • 如果使用 flattened_ids，則為 num_processes。
• (C4) 0 <= replica_groups < size(replica_groups)。
• (C5) 如果 use_global_device_ids = true，則 channel_id > 0。
• (C6) type(results...) = type(operands...) 除下列情況外：
  • dim(results..., all_gather_dim) = dim(operands..., all_gather_dim) * dim(process_groups, 1)。

範例

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

更多範例

all_reduce

語義學

在 StableHLO 程序網格中的每個程序群組內，將縮減函式 computation 套用至每個程序的 operands 張量值，並產生 results 張量。

這項作業會將 StableHLO 處理格線分割為 process_groups，其定義如下：

• cross_replica(replica_groups) 若 channel_id <= 0 and use_global_device_ids = false。
• 如果為 channel_id > 0 and use_global_device_ids = false，則為 cross_replica_and_partition(replica_groups)。
• flattened_ids(replica_groups) 若 channel_id > 0 and use_global_device_ids = true。

接著，在每個 process_group 中：

• 部分二進位檔樹狀結構 schedule 的 results...@process[result_index] = exec(schedule)，其中：
  • exec(node) = computation(exec(node.left), exec(node.right))。
  • exec(leaf) = leaf.value。
• schedule 是實作定義的二元樹，其內部順序遍歷是 to_destination_type(operands...@process_group...[result_index], type(func_inputs(computation)[0]))。

輸入

輸出內容

限制

• (C1) is_unique(replica_groups)。
• (C2) size(replica_groups) 的定義為：
  • 如果使用 cross_replica，則為 num_replicas。
  • 如果使用 cross_replica_and_partition，則為 num_replicas。
  • 如果使用 flattened_ids，則為 num_processes。
• (C3) 0 <= replica_groups < size(replica_groups)。
• (C4) 如果 use_global_device_ids = true，則 channel_id > 0。
• (C5) computation 具有 (tensor<E>, tensor<E>) -> (tensor<E>) 類型，其中 is_promotable(element_type(operand), E)。
• (C6) shape(results...) = shape(operands...)。
• (C7) element_type(results...) = E。

範例

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

更多範例

all_to_all

語義學

在 StableHLO 程序格線的每個程序群組中，沿著 split_dimension 將 operands 張量的值分割成部分，在程序之間散布分割的部分，沿著 concat_dimension 連結散布的部分，並產生 results 張量。這項作業會將 StableHLO 處理格線分割為 process_groups，其定義如下：

• 若 channel_id <= 0 則設為 cross_replica(replica_groups)。
• 若 channel_id > 0 則設為 cross_partition(replica_groups)。

之後，在每個 process_group 內：

• split_parts...@sender = split(operands...@sender, split_count, split_dimension) process_group 中的所有 sender。
• scattered_parts...@receiver = [split_parts...@sender[receiver_index] for sender in process_group]，其中 receiver_index = process_group.index(receiver)。
• results...@process = concatenate(scattered_parts...@process, concat_dimension)。

輸入

輸出內容

限制

• (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) 的定義為：
  • 如果使用 cross_replica，則為 num_replicas。
  • 如果使用 cross_partition，則為 num_partitions。
• (C7) 0 <= replica_groups < size(replica_groups)。
• (C8) dim(replica_groups, 1) = split_count。
• (C9) type(results...) = type(operands...)，但 split_dimension != concat_dimension 除外：
  • dim(results..., split_dimension) = dim(operands..., split_dimension) / split_count。
  • dim(results..., concat_dimension) = dim(operands..., concat_dimension) * split_count。

範例

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

更多範例

和

語義學

對兩個張量 lhs 和 rhs 執行元素 AND 運算，並產生 result 張量。根據元素類型執行以下操作：

• 布林值：邏輯 AND。
• 整數：位元 AND。

輸入

輸出內容

限制

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

範例

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

更多範例

atan2

語義學

對 lhs 和 rhs 張量執行元素級別 atan2 運算，並產生 result 張量。視元素類型而定，執行下列操作：

• 浮點值：IEEE-754 的 atan2。
• 複數：complex atan2。
• 對於量化類型：dequantize_op_quantize(atan2, lhs, rhs, type(result))。

輸入

輸出內容

限制

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

範例

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

更多範例

batch_norm_grad

語義學

計算 batch_norm_training 從 grad_output 反向傳播的多個輸入的梯度，並產生 grad_operand、grad_scale 和 grad_offset 張量。更正式地說，此作業可使用 Python 語法表示為現有的 StableHLO 作業進行分解，如下所示：

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

如果是量化類型，請執行 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))。

輸入

輸出內容

限制

• (C1) 0 <= feature_index < rank(operand)。
• (C2) operand、scale、mean、variance、grad_output、grad_operand、grad_scale 和 grad_offset 具有相同的 baseline_element_type。
• (C3) operand、grad_output 和 grad_operand 具有相同的形狀。
• (C4) scale、mean、variance、grad_scale 和 grad_offset 具有相同的形狀。
• (C5) size(scale) = dim(operand, feature_index)。

範例

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

語義學

會將 operand 張量在所有維度中進行規格化，但不包括 feature_index 維度，並產生 result 張量。更正式地說，這個運算可使用 Python 語法表示為對現有 StableHLO 運算的細分，如下所示：

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)

針對量化類型，執行 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))。

輸入

輸出內容

限制

• (C1) 0 <= feature_index < rank(operand)。
• (C2) operand、scale、offset、mean、variance 和 result 具有相同的 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)。

範例

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

語義學

計算所有維度的平均值和變異數 (feature_index 維度除外)，並將 operand 張量標準化，產生 output、batch_mean 和 batch_var 張量。更正式地說，這個運算可使用 Python 語法表達為現有 StableHLO 運算的解構，如下所示：

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

針對量化類型，執行 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))。

輸入

輸出內容

限制

• (C1) 0 <= feature_index < rank(operand)。
• (C2) operand、scale、offset、batch_mean、batch_var 和 output 具有相同的 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)。

範例

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

語義學

對 operand 張量執行位元轉換運算，並產生 result 張量，其中整個 operand 張量的位元會使用 result 張量的類型重新解讀。

更正式地說，假設有 E = element_type(operand)、E' = element_type(result) 和 R = rank(operand)：

• 如果是 num_bits(E') < num_bits(E)，則為 bits(result[i0, ..., iR-1, :]) = bits(operand[i0, ..., iR-1])。
• 如果是 num_bits(E') > num_bits(E)，則為 bits(result[i0, ..., iR-2]) = bits(operand[i0, ..., iR-2, :])。
• 如果是 num_bits(E') = num_bits(E)，則為 bits(result[i0, ..., iR-1]) = bits(operand[i0, ..., iR-1])。

bits 會傳回指定值的記憶體內表示法，而其行為是由實作定義，因為張量的確切表示法是由實作定義，元素類型的確切表示法也是由實作定義。

輸入

輸出內容

限制

• (C1) 指定 E = is_quantized(operand) ? storage_type(operand) : element_type(operand)、E' = is_quantized(result) ? storage_type(result) : element_type(result) 和 R = rank(operand)：
  • 如果是 num_bits(E') = num_bits(E)，shape(result) = shape(operand)。
  • 如果為 num_bits(E') < num_bits(E)：
  • rank(result) = R + 1。
  • 所有0 <= i < R的dim(result, i) = dim(operand, i)。
  • dim(result, R) * num_bits(E') = num_bits(E)。
  • 如果為 num_bits(E') > num_bits(E)：
  • rank(result) = R - 1。
  • dim(result, i) = dim(operand, i) 適用於所有 0 <= i < R。
  • dim(operand, R - 1) * num_bits(E) = num_bits(E')。
• (C2) 如果 is_complex(operand) or is_complex(result)，則為 is_complex(operand) and is_complex(result)。

範例

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

更多範例

broadcast_in_dim

語義學

透過複製 operand 張量中的資料，並產生 result 張量，擴充輸入張量的維度和/或秩。更正式的 result[result_index] = operand[operand_index]，其中 axes(operand) 中的所有 d：

• 若 dim(operand, d) = 1 則設為 operand_index[d] = 0。
• 否則傳回 operand_index[d] = result_index[broadcast_dimensions[d]]。

輸入

輸出內容

限制

• (C1) element_type(result) 的提供者：
  • element_type(operand)，如果 !is_per_axis_quantized(operand)。
  • element_type(operand) 但 quantization_dimension(operand)、scales(operand) 和 zero_points(operand) 可能與 quantization_dimension(result)、scales(result) 和 zero_points(result) 之間不同。
• (C2) size(broadcast_dimensions) = rank(operand)。
• (C3) 0 <= broadcast_dimensions < rank(result)。
• (C4) is_unique(broadcast_dimensions)。
• (C5) 針對 axes(operand) 中的所有 d：
  • dim(operand, d) = 1或
  • dim(operand, d) = dim(result, broadcast_dimensions[d])。
• (C6) 如果 is_per_axis_quantized(result)：
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)]。
  • 如果是 dim(operand, quantization_dimension(operand)) = 1，則 scales(result)[i] = scales(operand)[0] and zero_points(result)[i] = zero_points(operand)[0] for i in range(dim(result, quantization_dimension(result)))。

範例

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

更多範例

保護殼

語義學

根據 index 的值，從 branches 執行一項函式產生輸出內容。更正式的 result = selected_branch()，其中：

• 若 0 <= index < size(branches) 則設為 selected_branch = branches[index]。
• 否則傳回 selected_branch = branches[-1]。

輸入

輸出內容

限制

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

範例

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

更多範例

cbrt

語義學

對 operand 張量執行元素級別的立方根運算，並產生 result 張量。視元素類型而定，執行下列操作：

• 浮點值：IEEE-754 的 rootn(x, 3)。
• 複數：複數立方根。
• 針對量化類型：dequantize_op_quantize(cbrt, operand, type(result))

輸入

輸出內容

限制

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

範例

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

更多範例

ceil

語義學

對 operand 張量執行元素逐項的圓整，並產生 result 張量。實作 IEEE-754 規格中的 roundToIntegralTowardPositive 運算。如果是量化類型，請執行 dequantize_op_quantize(ceil, operand, type(result))。

輸入

輸出內容

限制

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

範例

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

更多範例

cholesky

語義學

計算一批矩陣的 Cholesky 分解。

更正式地說，對於 index_space(result) 中的所有 i，result[i0, ..., iR-3, :, :] 是 a[i0, ..., iR-3, :, :] 的 Cholesky 分解，以下三角 (如果 lower 是 true) 或上三角 (如果 lower 是 false) 矩陣的形式呈現。相反三角形中的輸出值 (即嚴格上三角形或嚴格下三角形) 則由實作定義。

如果輸入矩陣不是 Hermitian 正定矩陣的 i，則行為為未定義。

針對量化類型，執行 dequantize_op_quantize(lambda operand: cholesky(operand, lower), a, type(result))。

輸入

輸出內容

限制

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

範例

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

限制取值範圍

語義學

將 operand 張量的每個元素限制在最小值和最大值之間，並產生 result 張量。更正式的 result[result_index] = minimum(maximum(operand[result_index], min_element), max_element)，其中 min_element = rank(min) = 0 ? min[] : min[result_index]、max_element = rank(max) = 0 ? max[] : max[result_index]。針對量化類型，執行 dequantize_op_quantize(clamp, min, operand, max, type(result))。

對複數強制排序會涉及意料之外的語意，因此我們未來預計會移除此運算對複數的支援 (#560)。

輸入

輸出內容

限制

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

範例

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

更多範例

collective_broadcast

語義學

在 StableHLO 程序格線的每個程序群組中，將來源程序的 operand 張量值傳送至目標程序，並產生 result 張量。

該作業會將 StableHLO 程序格線分割為 process_groups，定義如下：

• 若 channel_id <= 0 則設為 cross_replica(replica_groups)。
• 如果 channel_id > 0，則為 cross_partition(replica_groups)。

之後，result@process 的值會由下列公式計算：

• 如果有 i 導致程序位於 process_groups[i]，則為 operand@process_groups[i, 0]。
• 否則為 broadcast_in_dim(constant(is_quantized(result) ? quantize(0, element_type(result)) : 0, element_type(result)), [], type(result))。

輸入

輸出內容

限制

• (C1) is_unique(replica_groups)。
• (C2) 0 <= replica_groups < N，其中 N 定義為：
  • 如果使用 cross_replica，則為 num_replicas。
  • 如果使用 cross_partition，則為 num_partitions。
• (C3) type(result) = type(operand)。

範例

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

語義學

在 StableHLO 程序格線的每個程序群組中，將 operand 張量的值從來源程序傳送至目標程序，並產生 result 張量。

這項作業會將 StableHLO 處理格線分割為 process_groups，其定義如下：

• 如果 channel_id <= 0，則為 cross_replica(source_target_pairs)。
• 如果 channel_id > 0，則為 cross_partition(source_target_pairs)。

之後，result@process 的值會由下列公式計算：

• operand@process_groups[i, 0]，如果存在 i，則 process_groups[i, 1] = process。
• 否則為 broadcast_in_dim(constant(is_quantized(result) ? quantize(0, element_type(result)) : 0, element_type(result)), [], type(result))。

輸入

輸出內容

限制

• (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，其中 N 的定義如下：
  • 如果使用 cross_replica，則為 num_replicas。
  • 如果使用 cross_partition，則為 num_partitions。
• (C5) type(result) = type(operand)。

範例

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

更多範例

比較

語義學

根據 comparison_direction 和 compare_type 執行 lhs 和 rhs 張元素依據元素的比較，並產生 result 張量。

comparison_direction 和 compare_type 的值具有以下語義：

布林值和整數元素類型：

• EQ：lhs = rhs。
• NE：lhs != rhs。
• GE：lhs >= rhs。
• GT：lhs > rhs。
• LE：lhs <= rhs。
• LT：lhs < rhs。

針對具有 compare_type = FLOAT 的浮點元素類型，運算會實作下列 IEEE-754 作業：

• EQ：compareQuietEqual。
• NE：compareQuietNotEqual。
• GE：compareQuietGreaterEqual。
• GT：compareQuietGreater。
• LE：compareQuietLessEqual。
• LT：compareQuietLess。

對於含有 compare_type = TOTALORDER 的浮點元素類型，運算子會使用 IEEE-754 中的 totalOrder 和 compareQuietEqual 運算組合。

針對複雜的元素類型，系統會使用提供的 comparison_direction 和 compare_type，執行 (real, imag) 組合的字典排序比對。對複數強制排序會涉及意義不明的語意，因此我們未來計劃在 comparison_direction 為 GE、GT、LE 或 LT 時，移除對複數的支援 (#560)。

針對量化類型，請執行 dequantize_compare(lhs, rhs, comparison_direction)。

輸入

輸出內容

限制

• (C1) baseline_element_type(lhs) = baseline_element_type(rhs)。
• (C2) shape(lhs) = shape(rhs) = shape(result)。
• (C3) compare_type 定義為：
  • 若 is_signed_integer(element_type(lhs)) 則設為 SIGNED。
  • 如果 is_unsigned_integer(element_type(lhs)) or is_boolean(element_type(lhs))，則為 UNSIGNED。
  • 如果是 is_float(element_type(lhs))，則為 FLOAT 或 TOTALORDER。
  • 若 is_complex(element_type(lhs)) 則設為 FLOAT。

範例

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

更多範例

複雜

語義學

這個外掛程式能從一組實際值和虛構值 (lhs 和 rhs) 執行元素轉換到複雜值，並產生 result 張量。

輸入

輸出內容

限制

• (C1) type(lhs) = type(rhs)。
• (C2) shape(result) = shape(lhs)。
• (C3) element_type(result) 具有 complex<E> 類型，其中 E = element_type(lhs)。

範例

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

更多範例

複合

語義學

封裝其他 StableHLO 作業 (由其他 StableHLO 作業組成) 的作業，執行 inputs 和 composite_attributes 並產生 results。decomposition 屬性會實作作業的語意。composite 運算子可以替換為其分解項目，而不會變更程式語意。如果內嵌分解作業無法提供相同的運算子語義，建議使用 custom_call。

version 欄位 (預設為 0) 用於表示組合的語意變更時間。

輸入

輸出內容

限制

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

範例

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

更多範例

concatenate

語義學

按照與指定引數相同的順序，沿著 dimension 維度連接 inputs，並產生 result 張量。更正式的說法是 result[i0, ..., id, ..., iR-1] = inputs[k][i0, ..., kd, ..., iR-1]，其中：

1. id = d0 + ... + dk-1 + kd。
2. d 等於 dimension，而 d0、... 是 inputs 的 d 維度大小。

輸入

輸出內容

限制

• (C1) same(element_type(inputs...))。
• (C2) same(shape(inputs...)) (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]) 除下列情況外：
  • dim(result, dimension) = dim(inputs[0], dimension) + ...。

範例

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

更多範例

常數

語義學

從常數 value 產生 output 張量。

輸入

輸出內容

限制

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

範例

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

更多範例

完成轉換

語義學

這個外掛程式能在 operand 張量執行元素類型之間的元素轉換，並產生 result 張量。

對於布林值至任何支援類型轉換，值 false 會轉換為零，而值 true 會轉換為一。對於any-supported-type-to-boolean的轉換，零值會轉換為 false，非零值會轉換為 true。請參閱下文，瞭解此做法如何適用於複雜類型。

如果轉換作業涉及整數轉整數、整數轉浮點或浮點轉浮點，如果來源值可在目的地類型中精確表示，結果值就是該精確表示法。否則，行為未定 (#180)。

針對涉及floating-point-to-integer的轉換，系統會截斷分數部分。如果截斷的值無法以目的地類型表示，行為為待定 (#180)。

涉及複雜至複雜的轉換會採用與浮點至浮點轉換相同的行為模式，轉換真實與虛構部分。

對於複雜至任何其他類型和任何其他類型至複雜轉換，系統會分別忽略來源虛構值或將目的地虛構值設為零。實部份的轉換會遵循浮點轉換。

原則上，這個運算可表示去量化 (從量化張量轉換為一般張量)、量化 (從一般張量轉換為量化張量) 和重新量化 (在量化張量之間轉換)，但目前我們有專屬的運算可用於此 - uniform_dequantize 適用於第一個用途，uniform_quantize 適用於第二個和第三個用途。日後，這兩個運算可能會合併為 convert (#1576)。

輸入

輸出內容

限制

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

範例

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

更多範例

卷積

語義學

計算 lhs 視窗與 rhs 切片之間的內積，並產生 result。下圖以具體範例說明如何從 lhs 和 rhs 計算 result 中的元素。

更正式地說，請考慮以下以 lhs 為準的輸入內容重構，以便表達 lhs 的視窗：

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

這項重構會使用下列輔助函式：

• 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] 其中 j[d] = i[permutation[d]]。

如果 feature_group_count = 1 和 batch_group_count = 1，則針對 index_space(dim(result, output_spatial_dimensions...)) 中的所有 output_spatial_index，result[result_shape(:, output_spatial_index, :)] = dot_product 為：

• 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])。這項功能似乎未使用，因此我們預計日後會移除這項功能 (#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])。

如果 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)。

如果 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)。

針對量化類型，執行 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))。

針對混合量化類型，執行 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)。

輸入

輸出內容

限制

• (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) 假設 input_dimensions = [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension]：
  • 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) Given kernel_dimensions = kernel_spatial_dimensions + [kernel_input_feature_dimension] + [kernel_output_feature_dimension]:
  • is_unique(kernel_dimensions)。
  • 0 <= kernel_dimensions < N。
• (C19) size(output_spatial_dimensions) = N - 2。
• (C20) 假設 output_dimensions = [output_batch_dimension] + output_spatial_dimensions + [output_feature_dimension]：
  • 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) 的定義為：
  • 若 result_dim = output_batch_dimension 則設為 dim(lhs, input_batch_dimension) / batch_group_count。
  • 如果 result_dim = output_feature_dimension，則為 dim(rhs, kernel_output_feature_dimension)。
  • num_windows 否則：
  • 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。
• 如果運算使用非量化的張量：
  • (C27) element_type(lhs) = element_type(rhs) = element_type(result)。
• 如果作業使用量化張量：
  • (C28) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs)。
  • (C29) 如果 is_per_axis_quantized(rhs)，則 quantization_dimension(rhs) = kernel_output_feature_dimension。
  • (C30) 如果是 is_per_axis_quantized(result)，則 quantization_dimension(result) = output_feature_dimension。
  • 如果 is_quantized(lhs)：
  • (C31) storage_type(lhs) = storage_type(rhs)。
  • (C32) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result)。
  • (C33) 如果 is_per_tensor_quantized(rhs)，則為 is_per_tensor_quantized(result)。
  • 如果 !is_quantized(lhs)：
  • (C34) element_type(lhs) = expressed_type(rhs) = element_type(result).

範例

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

更多範例

餘弦

語義學

對 operand 張量執行元素級別餘弦運算，並產生 result 張量。視元素類型而定，執行下列操作：

• 浮點值：IEEE-754 的 cos。
• 適用於複數：複數。
• 量化類型：dequantize_op_quantize(cosine, operand, type(result))。

輸入

輸出內容

限制

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

範例

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

更多範例

count_leading_zeros

語義學

按元素計算 operand 張前方的零位元數，並產生 result 張量。

輸入

輸出內容

限制

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

範例

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

更多範例

custom_call

語義學

封裝實作定義的作業 call_target_name，這會使用 inputs 和 called_computations 並產生 results。has_side_effect、backend_config 和 api_version 可用於提供其他實作定義的中繼資料。

目前，此運算包含相當雜亂的中繼資料集合，反映了 XLA 編譯器中對應運算的自然演進。我們計劃日後要統合這個中繼資料 (#741)。

輸入

輸出內容

範例

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

除號

語義學

對除數 lhs 和除數 rhs 張量執行元素逐元素除法，並產生 result 張量。視元素類型而定，執行下列操作：

• 整數：整數除法，產生代數商，並捨棄任何小數部分。
• 浮點值：IEEE-754 的 division。
• 複數：複數除法。
• 針對量化型別：
  • dequantize_op_quantize(divide, lhs, rhs, type(result))。

輸入

輸出內容

限制

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

範例

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

更多範例

dot_general

語義學

計算 lhs 切片和 rhs 切片之間的內積，並產生 result 張量。

更正式的說法是 result[result_index] = dot_product，其中：

• 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，其中 size(result_batching_index) = size(lhs_batching_dimensions)、size(result_lhs_index) = size(lhs_result_dimensions) 和 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))。

針對量化類型，執行 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))。

針對混合量化類型，執行 hybrid_dequantize_then_op( lambda lhs, rhs: dot_general(lhs, rhs, lhs_batching_dimensions, rhs_batching_dimensions, lhs_contracting_dimensions, rhs_contracting_dimensions, precision_config), lhs, rhs)。

precision_config 會控制加速器後端運算的速度與準確度之間的權衡。可以是下列其中一個值 (目前這些列舉值的語意未明確指定，但我們預計會在 #755 中解決這個問題)：

• DEFAULT：計算速度最快，但近似於原始數字的近似值不精確。
• HIGH：計算速度較慢，但可更準確地估算原始數字。
• HIGHEST：計算速度最慢，但最接近原始數字。

DotAlgorithm 會定義用於實作點運算法的必要屬性，並定義精確度。如果設定了演算法屬性欄位，則 precision_config 必須為 DEFAULT。DotAlgorithms 沒有預設值，因為預設參數是由實作定義。因此，所有點狀演算法欄位都可能設為 None，以指定空白點狀演算法，而該演算法會改用 precision_config 值。

DotAlgorithm 欄位包括：

• lhs_precision_type 和 rhs_precision_type，作業的 LHS 和 RHS 會四捨五入至此精確度。精確度類型與輸入和輸出內容的儲存類型無關。
• accumulation_type 用於累加的精確度。
• lhs_component_count、rhs_component_count 和 num_primitive_operations 的運算方法適用於將 LHS 和/或 RHS 分解為多個元件，並對這些值執行多項「原始」點運算時，通常用於模擬較高的精確度 (例如運用 bfloat16 bfloat16 picial Intelligence Datatype for Higher-Precision tf2put6_6：如果是沒有分解的演算法，則應將這些值設為 1。
• allow_imprecise_accumulation 可用於指定是否允許在某些步驟中使用較低精確度的累加運算 (例如 CUBLASLT_MATMUL_DESC_FAST_ACCUM)。

DotAlgorithm 屬性範例：

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


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


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

實作方式會決定支援哪些組合。一般而言，無法保證 StableHLO 消費者支援各種加速器類型。如果不支援特定演算法，則應引發錯誤，而非改回使用替代方法。StableHLO 驗證會盡力驗證，避免在任何硬體上支援未知的演算法。

如要瞭解部分支援的演算法值，請參閱 xla_data.proto > Algorithm。支援單 #2483 記錄了建立一份關於後端支援演算法的集中式文件的計畫。

輸入

輸出內容

限制

• (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)。
• 如果作業使用非量化張量：
  • (C13) element_type(lhs) = element_type(rhs)。
• 如果運算使用經過量化的張量：
  • (C14) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs)。
  • (C15) zero_points(rhs) = 0。
  • (C16) 如果 is_per_axis_quantized(rhs)，則 quantization_dimension(rhs) 不在 rhs_contracting_dimensions 中。
  • 如果 is_quantized(lhs)：
  • (C17) storage_type(lhs) = storage_type(rhs)。
  • (C18) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result)。
  • (C19) 如果 is_per_tensor_quantized(rhs)，則為 is_per_tensor_quantized(result)。
  • 如果 !is_quantized(lhs)：
  • (C20) element_type(lhs) = expressed_type(rhs) = element_type(result)。
• 如果 !is_empty_algorithm(lhs_precision_type, rhs_precision_type, accumulation_type, lhs_component_count, rhs_component_count, num_primitive_operations allow_imprecise_accumulation)：
  • (C21) precision_config... = DEFAULT。
  • (C22) 0 < lhs_component_count。
  • (C23) 0 < rhs_component_count。
  • (C24) 0 < num_primitive_operations。

範例

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

更多範例

dynamic_broadcast_in_dim

語義學

這項作業的功能與 broadcast_in_dim 作業相同，但結果形狀會透過 output_dimensions 動態指定。

這項作業也接受選用屬性 known_expanding_dimensions、known_nonexpanding_dimensions，用來表示有關維度展開行為的靜態知識。如果未指定，系統會假設所有維度都可能會展開。

輸入

輸出內容

限制

• (C1) element_type(result) 的提供者：
  • element_type(operand)，如果 !is_per_axis_quantized(operand)。
  • element_type(operand) 但 quantization_dimension(operand)、scales(operand) 和 zero_points(operand) 可能與 quantization_dimension(result)、scales(result) 和 zero_points(result) 之間不同。
• (C2) size(broadcast_dimensions) = rank(operand)。
• (C3) 0 <= broadcast_dimensions < rank(result)。
• (C4) is_unique(broadcast_dimensions)。
• (C5) 針對 axes(operand) 中的所有 d：
  • dim(operand, d) = 1或
  • dim(operand, d) = dim(result, broadcast_dimensions[d])。
• (C6) 如果 is_per_axis_quantized(result)：
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)]。
  • 如果是 dim(operand, quantization_dimension(operand)) = 1，則 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_nonexpanding_dimensions)。
• (C9) 0 <= known_expanding_dimensions < rank(operand)。
• (C10) 0 <= known_nonexpanding_dimensions < rank(operand)。

範例

// %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_nonexpanding_dimensions = array<i64: 1>
} : (tensor<1x3xi64>, tensor<3xi64>) -> tensor<2x3x2xi64>
// %result: [
//            [
//             [1, 1],
//             [2, 2],
//             [3, 3]
//            ],
//            [
//             [1, 1],
//             [2, 2],
//             [3, 3]
//            ]
//          ]

更多範例

dynamic_conv

語義學

這項作業的功能與卷積作業相同，但邊框會透過 padding 動態指定。

輸入

輸出內容

限制

• (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) 假設 input_dimensions = [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension]：
  • 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) Given kernel_dimensions = kernel_spatial_dimensions + [kernel_input_feature_dimension] + [kernel_output_feature_dimension]:
  • is_unique(kernel_dimensions)。
  • 0 <= kernel_dimensions < N。
• (C19) size(output_spatial_dimensions) = N - 2。
• (C20) 假設 output_dimensions = [output_batch_dimension] + output_spatial_dimensions + [output_feature_dimension]：
  • 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) 的定義為：
  • 若 result_dim = output_batch_dimension 則設為 dim(lhs, input_batch_dimension) / batch_group_count。
  • 如果 result_dim = output_feature_dimension，則為 dim(rhs, kernel_output_feature_dimension)。
  • num_windows 否則：
  • 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。
• 如果運算使用非量化的張量：
  • (C27) element_type(lhs) = element_type(rhs) = element_type(result)。
• 如果作業使用量化張量：
  • (C28) is_quantized(lhs) = is_quantized(result) and is_quantized(rhs)。
  • (C29) 如果 is_per_axis_quantized(rhs)，則 quantization_dimension(rhs) = kernel_output_feature_dimension。
  • (C30) 如果是 is_per_axis_quantized(result)，則 quantization_dimension(result) = output_feature_dimension。
  • 如果 is_quantized(lhs)：
  • (C31) storage_type(lhs) = storage_type(rhs)。
  • (C32) expressed_type(lhs) = expressed_type(rhs) = expressed_type(result)。
  • (C33) 如果 is_per_tensor_quantized(rhs)，則為 is_per_tensor_quantized(result)。
  • 如果 !is_quantized(lhs)：
  • (C34) element_type(lhs) = expressed_type(rhs) = element_type(result).

範例

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