StableHLO の仕様

StableHLO は、マシンでの高レベル操作（HLO）のためのオペレーション セット 学習（ML）モデルです。StableHLO は、異なるサーバー間のポータビリティ レイヤとして機能します。 ML フレームワークと ML コンパイラ: StableHLO プログラムを生成する ML フレームワーク StableHLO プログラムを使用する ML コンパイラと互換性がある

Google の目標は、より多くのプロダクトを開発することで、ML 開発を簡素化、加速させることです。 さまざまな ML フレームワーク（TensorFlow、JAX、 PyTorch など）や ML コンパイラ（XLA や IREE など）を使用したデータ処理パイプラインの構築をサポートします。そのために、 に、StableHLO プログラミング言語の仕様が記載されています。

この仕様には、大きく分けて 3 つのセクションがあります。まず、 [プログラム] セクションでは、StableHLO プログラムの構造について説明します。 これは StableHLO オペレーションで構成され、それ自体が StableHLO オペレーションで構成されます。 その構造内の Ops セクションでは、 見ていきますExecution セクションでは、すべての 1 つのプログラム内で 一緒に実行されるからです最後に、 表記セクションでは、本コース全体で使用される表記法を 仕様。

StableHLO の以前のリリースの仕様を表示するには、 タグ付けされたリリースです。 たとえば、StableHLO v0.19.0 仕様などです。 StableHLO のマイナー バージョンのバンプごとに発生した変更を表示するには、以下をご覧ください。 VhloDialect.td にあるバージョン ログ。

プログラム

Program ::= {Func}

StableHLO プログラムは、任意の数の StableHLO 関数で構成されます。 以下は、3 つの入力を持つ関数 @main を使用したプログラムの例を示しています。 （%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 関数（名前付き関数とも呼ばれます）には、 識別子、入出力、本文があります将来的には 互換性を高めるために関数のメタデータを追加導入する （#425、 #626、 #740、 #744）。

識別子

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

StableHLO 識別子は、多くのプログラミングの識別子に似ています。 2 つの特徴があります。1）すべての識別子には、 異なる種類の識別子を区別する、2）値の識別子を 完全に数値化して 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} | '?'

テンソル型はテンソル、つまり多次元配列を表します。この会社には shape と element type: シェイプは負でないか、 未知のディメンション サイズ（対応する要素の昇順） ディメンション（軸とも呼ばれます）は 0 から R-1 まで番号が付けられます。「 次元数 R はランクと呼ばれます。たとえば、tensor<2x3xf32> は次のようになります。 形状が 2x3、要素型が f32 のテンソル型。2 つの次元があります。 （つまり、2 つの軸）- 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 ::= IntegerConstant
QuantizationStorageMax ::= IntegerConstant
QuantizationExpressedType ::= FloatType
QuantizationDimension ::= IntegerConstant
QuantizationParameters ::= QuantizationParameter
                         | '{' QuantizationParameter {',' QuantizationParameter} '}'
QuantizationParameter ::= QuantizationScale ':' QuantizationZeroPoint
QuantizationScale ::= FloatConstant
QuantizationZeroPoint ::= IntegerConstant

量子化要素型は、ストレージ タイプの整数値を表します。 対応する 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 のセマンティクスについては、現在議論が続いています。 型、値のほか、1 つまたは 1 つの指標が ゼロ点が複数含まれることがあります。基づいて、 0 ポイント前後の仕様は変わる可能性があります。 （#1405）。

QuantizationStorageMin のセマンティクスに関して進行中の別のディスカッションがあります と QuantizationStorageMax を使用して、制約を適用するかどうかを決定します。 これらの値と量子化テンソルの値に適用される （#1406）。

最後に、未知のスケールとゼロの表現について 未知の表現をモデル化する方法について （#1407）

量子化テンソル型は、量子化された要素を持つテンソルを表します。これらの テンソルは通常のテンソルとまったく同じですが、テンソルの要素が 通常の要素型ではなく、量子化された要素型を使用します。

量子化テンソルでは、量子化はテンソルごとに行うことができます。 テンソル全体で 1 つの scale と zero_point、または軸ごとにすることもできる つまり、複数の scales と zero_points があり、スライスごとに 1 つのペアが quantization_dimension を返します。より正式には、テンソル t で表します。 軸ごとの量子化の場合、dim(t, quantization_dimension) 個のスライスがあります。 /quantization_dimension: t[:, ..., 0, ..., :], t[:, ..., 1, ..., :]、 i 番目のスライス内のすべての要素は、次のように scales[i] と zero_points[i] を使用します。 生成します。量子化テンソル型は次のとおりです。 制約:

• テンソルごとの量子化の場合: <ph type="x-smartling-placeholder">
  </ph>
  • 追加の制約はありません。
• 軸ごとの量子化の場合: <ph type="x-smartling-placeholder">
  </ph>
  • （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 ::= 'f8E4M3FN' | 'f8E5M2' | 'f8E4M3FNUZ' | 'f8E5M2FNUZ'
            | 'f8E4M3B11FNUZ' | 'bf16' | 'f16' | 'f32' | 'f64'
TensorFloat32 ::= 'tf32'
ComplexType ::= 'complex' '<' ComplexElementType '>'
ComplexElementType ::= 'f32' | 'f64'

要素型は、テンソル型の要素を表します。多くのプログラミングでは、 これらの型は StableHLO のファースト クラスではありません。つまり StableHLO プログラムは、これらの型の値を直接表すことはできません（そのため、 T 型のスカラー値を 0 次元のテンソルで表すのが慣用的です。 tensor<T> 型の値など）。

• ブール値型は、ブール値 true と false を表します。
• 整数型は、符号付き（si）または符号なし（ui）のいずれかです。 サポートされているビット幅（2、4、8、16、32、64）のいずれか。 符号付き siN 型は -2^(N-1) から 2^(N-1)-1 までの整数値を表す 符号なし uiN 型は、0 から uiN までの整数値を表します。 2^N-1 を含む。
• 浮動小数点型は次のいずれかになります。 <ph type="x-smartling-placeholder">
  </ph>
  • f8E4M3FN 型と f8E5M2 型はそれぞれ FP8 形式の E4M3 および E5M2 エンコード（ 「FP8 Formats for Deep Learning」。
  • E4M3 と E5M2 に対応する f8E4M3FNUZ 型と f8E5M2FNUZ 型 FP8 形式の 2 つのエンコードをサポートしており、 ディープ ニューラル ネットワーク用の 8 ビット数値形式。
  • FP8 形式の E4M3 エンコードに対応する f8E4M3B11FNUZ 型 説明 ディープ ニューラル ネットワークのためのハイブリッド 8 ビット浮動小数点（HFP8）トレーニングと推論。
  • 前述の bfloat16 形式に対応する bf16 型 BFloat16: Cloud TPU で高いパフォーマンスを実現する秘密。
  • それぞれに対応する f16、f32、f64 型。 binary16（「半精度」）、binary32（「単精度」）、 binary64（「倍精度」）形式（ IEEE 754 規格に準拠しています。
  • tf32 型は TensorFloat32 形式に対応します。 StableHLO でのサポートは 限定的です
• 複合型は、実数部を持つ複合値を表します。 同じ要素タイプの虚数部サポートされている複合型 型は complex<f32>（どちらも f32 型）と complex<f64> です。 （どちらの部分も f64 型です）。

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

関数型は、名前付き関数と匿名関数の両方を表します。持っている 入力型（-> の左側の型のリスト）と出力型 （-> の右側にある型のリスト）。多くのプログラミングでは、 関数型はファースト クラスですが、StableHLO ではそうではありません。

StringType ::= 'string'

文字列型はバイトのシーケンスを表します。多くのプログラミングでは、 文字列型は StableHLO のファースト クラスではなく、次の目的でのみ使用される プログラム要素の静的メタデータを指定します。

運用

StableHLO オペレーション（オペレーションとも呼ばれます）はクローズド セットを表す おおまかに説明しましたすでに説明したように StableHLO の構文は MLIR に大きく影響しますが、MLIR は 人間工学に基づいた代替手段ですが、StableHLO の目標である ML フレームワークと ML コンパイラ間の 相互運用性を向上できます

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

StableHLO オペレーション（オペレーションとも呼ばれます）には名前があります。 署名で構成されています。名前は、stablehlo. 接頭辞と サポートされているいずれかのオペレーションを一意に識別するニーモニック。下記をご覧ください: サポートされているすべてのオペレーションの 包括的なリストをご覧ください

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

Ops は入力を消費し、出力を生成します。入力は以下のように分類されます。 入力値（実行中に計算される）、入力関数（ StableHLO では関数はファースト クラス値ではないため、静的に保持されます。 入力属性（静的にも提供されます）入力と出力の種類 演算で消費、生成される情報はそのニモニックに左右されます。たとえば、add op は 2 つの入力値を消費し、1 つの出力値を生成します。これに対し、 select_and_scatter 演算は 3 つの入力値、2 つの入力関数、 3 つの入力属性。

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

入力関数（匿名関数とも呼ばれます）は、 名前付き関数に似ていますが、次の点が異なります。1）識別子がない（つまり 「anonymous」という名前）、2）出力タイプを宣言しない（出力タイプは 関数内の return オペレーションから推論されます）。

入力関数の構文には、現在使用されていない部分が含まれます（ 上記の 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 opset に吸収するか、 StableHLO プログラムに 表示されますそれまでは 属性:

• layout（#629）。
• mhlo.frontend_attributes （#628）。
• mhlo.sharding（#619）。
• output_operand_aliases （#740）。
• 位置情報のメタデータ（#594）。

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

演算シグネチャは、すべての入力値の型（ -> の左側）とすべての出力値の型（すべての出力値の -> の右側にあります）。厳密に言えば、入力タイプは 出力型もほぼ常に冗長です（ 出力の型は入力から推測できます）。それでも シグネチャは、MLIR との互換性を確保するために、StableHLO 構文に意図的に含まれています。

ニモニックが select_and_scatter の演算の例を次に示します。消費電力は 入力値（%operand、%source、%init_value）、2 つの入力関数 と 3 つの入力属性（window_dimensions、window_strides、padding）。 op のシグネチャには入力値の型のみが含まれることに注目 （ただし、インラインで提供される入力関数と属性のタイプは対象外です）。

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

整数定数は、10 進数または 16 進数表記で表します。その他のベース（例:2 進数または 8 進数はサポートされません。 整数定数には次の制約があります。

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

浮動小数点定数は、 小数または科学的表記を使用できますさらに、16 進数表記を を使用して、基になるビットを浮動小数点形式で直接指定します。 対応する種類です浮動小数点定数には次の制約があります。

• （C1）16 進数以外の表記が使われている場合 is_wellformed(float_literal, float_type)。
• （C2）16 進数表記の場合 size(hexadecimal_digits) = num_bits(float_type) / 4。

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

複素定数は実数部のリストを使用して複素数を表す （最初に来る）と虚数部（2 番目）があります。たとえば (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

テンソル定数は、 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))。各要素の意味は次のとおりです。 <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))。各要素の意味は次のとおりです。 <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:])。
  • それ以外の場合は 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 です。

運用

abs

セマンティクス

operand テンソルに対して要素単位の abs 演算を実行し、result を生成します。 テンソルです。要素のタイプに応じて、次の操作を行います。

• 符号付き整数の場合: 整数のモジュラス。
• 浮動小数点数: IEEE-754 の abs。
• 複素数の場合: 複素率。
• 量子化型の場合: dequantize_op_quantize(abs, operand, type(result))。

入力

出力

制約

• （C1）shape(result) = shape(operand)。
• （C2）baseline_element_type(result) は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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]

その他の例

追加

セマンティクス

2 つのテンソル lhs と rhs を要素ごとに加算し、 result のテンソル。要素のタイプに応じて、次の操作を行います。

• ブール値の場合: 論理 OR。
• 整数の場合: 整数の加算。
• 浮動小数点数: IEEE-754 の addition。
• 複素数の場合: 複素加算。
• 量子化型の場合: dequantize_op_quantize(add, lhs, rhs, type(result))。

入力

出力

制約

• オペレーションで量子化されていないテンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （C1）type(lhs) = type(rhs) = type(result)。
• オペレーションで量子化テンソルを使用する場合: <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）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 の場合。
• cross_replica_and_partition(replica_groups) channel_id > 0 and use_global_device_ids = false の場合。
• flattened_ids(replica_groups) channel_id > 0 and use_global_device_ids = true の場合。

その後、各 process_group 内で以下を行います。

• operands...@receiver = [operand@sender for sender in process_group]（すべて） receiver（process_group）。
• results...@process = concatenate(operands...@process, all_gather_dim)（すべて） process（process_group）。

入力

出力

制約

• （C1）0 <= all_gather_dim < rank(operands...)。
• （C2）is_unique(replica_groups)。
• （C3）size(replica_groups) は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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...)、ただし: <ph type="x-smartling-placeholder">
  </ph>
  • 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 の場合。
• cross_replica_and_partition(replica_groups) channel_id > 0 and use_global_device_ids = false の場合。
• flattened_ids(replica_groups) channel_id > 0 and use_global_device_ids = true の場合。

その後、各 process_group 内で以下を行います。

• results...@process[result_index] = exec(schedule)（バイナリツリーの場合） schedule 各要素の意味は次のとおりです。 <ph type="x-smartling-placeholder">
  </ph>
  • 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) は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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 プロセス グリッドの各プロセス グループ内で、 operands テンソルを split_dimension に沿って分割し、分割を分散させる 分散された部分を連結して、プロセス間の 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) は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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 の場合を除く）: <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。

例

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

その他の例

と

セマンティクス

2 つのテンソル 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。
• 複素数の場合: 複素数 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 が生成されます テンソルです。より形式的には、このオペレーションは、 既存の StableHLO オペレーションを次のように Python 構文で記述します。

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 テンソルを生成します。より正式には オペレーションは既存の StableHLO オペレーションへの分解として表現できる 次のように Python 構文を使用します。

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) がある場合: <ph type="x-smartling-placeholder">
  </ph>
  • 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。
  • すべての 0 <= i < R に対して dim(result, i) = dim(operand, i)。
  • 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] のすべての d axes(operand):

• dim(operand, d) = 1 の場合は operand_index[d] = 0。
• そうでない場合は operand_index[d] = result_index[broadcast_dimensions[d]]。

入力

出力

制約

• （C1）element_type(result) は次式で求められます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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 について: <ph type="x-smartling-placeholder">
  </ph>
  • dim(operand, d) = 1 または
  • dim(operand, d) = dim(result, broadcast_dimensions[d])。
• （C6）is_per_axis_quantized(result) の場合: <ph type="x-smartling-placeholder">
  </ph>
  • 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]
//            ]
//          ]

その他の例

ケース

セマンティクス

branches から関数を 1 つだけ実行すると、出力が生成されます index の値に応じて異なります。よりフォーマルな表現で、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 テンソルの要素単位の ceil を実行し、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, :, :]（下向きの三角形または （lower が true の場合）または上三角形（lower が false の場合）の行列。 反対側の三角形（厳密な上三角形）の出力値 それに対応して厳密な下三角形が定義され、実装定義になります。

入力行列がエルミート正定値でない 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 は次で求められます。

• operand@process_groups[i, 0]（プロセスが次のようになっているような i が存在する場合） （process_groups[i]）
• 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 は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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 は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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]]

その他の例

比較

セマンティクス

次に従って lhs テンソルと rhs テンソルの要素単位の比較を行います。 comparison_direction と compare_type があり、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 を持つ浮動小数点要素型の場合、演算は 次の totalOrder オペレーションと compareQuietEqual オペレーションを組み合わせて使用します。 IEEE-754

複雑な要素型の場合、(real, imag) ペアの辞書順比較は次のようになります。 （指定された comparison_direction と compare_type を使用して実行される） 複素数に順序を付けるには、意外な意味論が必要です。 将来的には複素数のサポートは終了し 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 は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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 オペレーションで構成された（構成）オペレーションをカプセル化します。 inputs と composite_attributes を取り、results を生成します。「 op のセマンティクスは、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 は d 番目のディメンション サイズです /inputs。

入力

出力

制約

• （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])、ただし以下を除く: <ph type="x-smartling-placeholder">
  </ph>
  • 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 テンソルを生成します。

boolean-to-any-supported-typeの場合、値 false は次のようになります。 0 に変換され、値 true は 1 に変換されます。対象 any-supported-type-to-booleanが行われる場合、ゼロ値は false とゼロ以外の値は true に変換されます。この仕組みについては、下記をご覧ください。 複雑な型で機能します。

整数から整数または整数から浮動小数点数を含む変換の場合 または floating-point-to-floating-point（ソース値が が宛先タイプで表されている場合、結果の値は 必要があります。それ以外の場合の動作は未定 （#180）。

floating-point-to-integerへの変換の場合、小数部は次のようになります。 切り捨てられます。切り捨てられた値を宛先タイプで表現できない場合は、 動作は未定です（#180）。

複合型を含む変換は、 floating-point-to-floating-pointの変換（実数と浮動小数点数） 使用します。

complex-to-any-other-type と any-other-type-to-complex の変換の場合: ソースの虚偽値が無視されるか、宛先の虚数が それぞれゼロに設定されます。実数部分の変換は、 浮動小数点変換も実行できます

原理的には、この演算は逆量子化（ 量子化テンソルから規則テンソルへの変換）、量子化（通常のテンソルから テンソルから量子化テンソルへの変換）と再量子化（量子化テンソルと テンソルがありますが、現時点では専用のオペレーションがあります。 最初のユースケースは uniform_dequantize、ユースケースは uniform_quantize です。 2 つ目と 3 つ目のユースケースです将来的には、この 2 つのオペレーションが統合される可能性があります。 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。次の図は、result の要素がどのように計算されるかを示しています。 lhs と rhs を具体例で説明します。

より正式には、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 の場合、すべての output_spatial_index（index_space(dim(result, output_spatial_dimensions...))）、 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] の場合: <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）ある 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）ある 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) は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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。
• オペレーションで量子化されていないテンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （C27）element_type(lhs) = element_type(rhs) = element_type(result)。
• オペレーションで量子化テンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （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 内の先頭 0 ビットの数を要素ごとにカウントします 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。
• 複素数の場合: 複素除算。
• 量子化型の場合: <ph type="x-smartling-placeholder">
  </ph>
  • 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 は、次の要素の速度と精度のトレードオフを制御します。 コンピューティングの負荷を分散します以下のいずれかを指定できます（ これらの列挙値のセマンティクスは指定されていませんが、 2024 年には #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 AI データ型を活用して高精度の計算を行う: bf16_6x tf32_3x など）。分解を行わないアルゴリズムの場合、これらの値は 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 の検証ではベスト エフォート型の検証、 どのハードウェアでもサポートされないことがわかっているアルゴリズムを防止できます。

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)。
• オペレーションで量子化されていないテンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （C13）element_type(lhs) = element_type(rhs)。
• オペレーションで量子化テンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （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) の場合: <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。

例

// %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_non_expanding_dimensions も使用できます。 ディメンションの展開動作に関する静的な知識を表すために使用します。 指定しない場合、すべてのディメンションが拡張される可能性があるものと見なされます。

入力

出力

制約

• （C1）element_type(result) は次式で求められます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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 について: <ph type="x-smartling-placeholder">
  </ph>
  • dim(operand, d) = 1 または
  • dim(operand, d) = dim(result, broadcast_dimensions[d])。
• （C6）is_per_axis_quantized(result) の場合: <ph type="x-smartling-placeholder">
  </ph>
  • 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_non_expanding_dimensions)。
• （C9）0 <= known_expanding_dimensions < rank(operand)。
• （C10）0 <= known_non_expanding_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_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]
//            ]
//          ]

その他の例

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] の場合: <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）ある 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）ある 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) は次のように定義されます。 <ph type="x-smartling-placeholder">
  </ph>
  • 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。
• オペレーションで量子化されていないテンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （C27）element_type(lhs) = element_type(rhs) = element_type(result)。
• オペレーションで量子化テンソルを使用する場合: <ph type="x-smartling-placeholder">
  </ph>
  • （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]]
//          ]]