StableHLO की खास बातें

StableHLO, मशीन में हाई-लेवल ऑपरेशन (HLO) के लिए इस्तेमाल होने वाला एक ऑपरेशन सेट है लर्निंग (एमएल) मॉडल. StableHLO, अलग-अलग प्लैटफ़ॉर्म के बीच पोर्टेबिलिटी लेयर के तौर पर काम करता है एमएल फ़्रेमवर्क और एमएल कंपाइलर: एमएल फ़्रेमवर्क जो StableHLO प्रोग्राम बनाते हैं एमएल कंपाइलर के साथ काम करते हैं, जो StableHLO प्रोग्राम का इस्तेमाल करते हैं.

हमारा मकसद, मशीन लर्निंग को आसान बनाने के साथ-साथ मशीन लर्निंग को बेहतर बनाने के लिए, मशीन लर्निंग सिस्टम को अलग-अलग एमएल फ़्रेमवर्क के बीच इंटरऑपरेबिलिटी (दूसरे सिस्टम के साथ काम करना), जैसे कि TensorFlow, JAX, और PyTorch) और ML कंपाइलर (जैसे कि XLA और IREE). उस ओर, यह इस दस्तावेज़ में StableHLO प्रोग्रामिंग भाषा की जानकारी दी गई है.

इस स्पेसिफ़िकेशन में तीन मुख्य सेक्शन शामिल हैं. सबसे पहले, प्रोग्राम सेक्शन में, StableHLO प्रोग्राम की बनावट के बारे में जानकारी दी जाती है इनमें StableHLO फ़ंक्शन शामिल होते हैं, जिनमें StableHLO ops शामिल होता है. उस स्ट्रक्चर में, Ops सेक्शन, व्यक्तिगत ऑपरेशन एक्ज़िक्यूशन सेक्शन सभी के लिए सिमेंटिक उपलब्ध कराता है ये Ops, किसी प्रोग्राम में एक साथ एक्ज़ीक्यूट किए जाते हैं. आख़िर में, नोटेशन सेक्शन में स्पेसिफ़िकेशन.

StableHLO की पिछली रिलीज़ की खास जानकारी देखने के लिए, रेपो खोलें पसंद की टैग की गई रिलीज़. उदाहरण के लिए, StableHLO v0.19.0 स्पेसिफ़िकेशन. StableHLO के हर छोटे वर्शन बंप पर होने वाले बदलावों को देखने के लिए, इसे देखें VhloDialect.td वर्शन में लॉग इन करें.

प्रोग्राम

Program ::= {Func}

StableHLO प्रोग्राम में आर्बिट्रेरी संख्या में StableHLO फ़ंक्शन होते हैं. यहां @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 फ़ंक्शन (जिन्हें नाम वाले फ़ंक्शन भी कहा जाता है) में एक आइडेंटिफ़ायर, इनपुट/आउटपुट, और एक बॉडी. आने वाले समय में, हम फ़ंक्शन के लिए अतिरिक्त मेटाडेटा की जानकारी दें, ताकि वे बेहतर तरीके से काम कर सकें HLO (#425, #626, #740, #744).

आइडेंटिफ़ायर

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

StableHLO आइडेंटिफ़ायर, कई प्रोग्राम में मौजूद आइडेंटिफ़ायर से मिलते-जुलते होते हैं होती है, जिसकी दो विशेषताएं होती हैं: 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} | '?'

टेन्सर के टाइप टेंसर यानी कई डाइमेंशन वाले सरणियों को दिखाते हैं. उनके पास आकार और एलिमेंट का टाइप, जिसमें आकार गैर-ऋणात्मक या संबंधित डाइमेंशन साइज़ के बढ़ते क्रम में डाइमेंशन (जिन्हें ऐक्सिस भी कहा जाता है) को 0 से R-1 तक के नंबर के साथ लिखा जाता है. कॉन्टेंट बनाने डाइमेंशन की संख्या R को रैंक कहा जाता है. उदाहरण के लिए, tensor<2x3xf32> 2x3 आकार और f32 एलिमेंट वाला टेंसर टाइप. इसमें दो डाइमेंशन हैं (या, दूसरे शब्दों में, दो ऐक्सिस) - पहला डाइमेंशन और पहला डाइमेंशन - जिसके साइज़ नंबर 2 और 3 हैं. इसकी रैंक दो है.

आकार आंशिक या पूरी तरह से अज्ञात (डाइनैमिक) हो सकता है, जैसे 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 के सिमैंटिक पर एक बातचीत चल रही है, इसमें वे टाइप और वैल्यू शामिल होती हैं. साथ ही, यह भी जानकारी होती है कि क्या सिर्फ़ एक या क्वांटाइज़्ड टेंसर टाइप में, कई शून्य पॉइंट हो सकते हैं. इसके आधार पर इस चर्चा के परिणाम, शून्य बिंदु के लिए निर्देश बदल सकते हैं आने वाले समय में (#1405).

एक अन्य चर्चा में QuantizationStorageMin के सिमेंटिक्स शामिल हैं और QuantizationStorageMax तय करते हैं कि किसी कंस्ट्रेंट को इन वैल्यू और क्वांटाइज़्ड टेंसर की वैल्यू पर लगाया गया (#1406).

आखिर में, हम अज्ञात स्केल और शून्य को दर्शाने वाले तरीके के बारे में जानने की योजना बना रहे हैं जैसे हम अनजान लोगों को दिखाने के लिए डाइमेंशन का साइज़ (#1407).

क्वांटाइज़्ड टेंसर टाइप, क्वांटाइज़्ड एलिमेंट वाले टेंसर दिखाते हैं. ये टेंसर, सामान्य टेंसर की तरह ही होते हैं, अंतर सिर्फ़ इतना होता है कि उनके एलिमेंट इसमें सामान्य एलिमेंट के बजाय क्वांटाइज़्ड एलिमेंट टाइप होते हैं.

क्वांटाइज़्ड टेंसर में, क्वांटाइज़ेशन हर टेंसर के हिसाब से हो सकता है. इसका मतलब है कि पूरे टेंसर के लिए एक scale और zero_point या हर ऐक्सिस पर एक वैल्यू हो सकती है, मतलब कई scales और zero_points होने का मतलब है, इमेज के हर स्लाइस के लिए एक जोड़ा कोई विशेष आयाम quantization_dimension. औपचारिक तौर पर, टेंसर में t हर एक्सिस पर संख्या के हिसाब से, dim(t, quantization_dimension) स्लाइस हैं 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>> किसी खास डेटा टाइप के आधार पर असल में अलग हो सकता है लागू करना. आने वाले समय में, हम एचएलओ एबीआई में बदलाव करने की योजना बना रहे हैं जिसकी मदद से हम 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 से लेकर इंटिजर वैल्यू को दिखाते हैं 2^N-1 के साथ.
• फ़्लोटिंग-पॉइंट के टाइप इनमें से कोई एक हो सकते हैं:
  • f8E4M3FN और f8E5M2 टाइप, जो FP8 फ़ॉर्मैट की E4M3 और E5M2 एन्कोडिंग में बताया गया है डीप लर्निंग के लिए FP8 फ़ॉर्मैट.
  • E4M3 और E5M2 से जुड़े f8E4M3FNUZ और f8E5M2FNUZ टाइप FP8 फ़ॉर्मैट की एन्कोडिंग में बताए गए हैं डीप न्यूरल नेटवर्क के लिए 8-बिट नंबर वाले फ़ॉर्मैट.
  • FP8 फ़ॉर्मैट की E4M3 एन्कोडिंग से जुड़ा f8E4M3B11FNUZ टाइप इसमें बताया गया है डीप न्यूरल नेटवर्क के लिए, हाइब्रिड 8-बिट फ़्लोटिंग पॉइंट (HFP8) ट्रेनिंग और अनुमान.
  • bf16 टाइप, जो इसमें बताए गए bfloat16 फ़ॉर्मैट से जुड़ा है BFloat16: Cloud TPU पर बेहतर परफ़ॉर्मेंस पाने का सीक्रेट.
  • इसी तरह के f16, f32, और f64 टाइप binary16 ("हाफ़ सटीक"), binary32 ("एकल सटीक") और binary64 ("दोगुने सटीक") फ़ॉर्मैट में बताया गया है आईईईई 754 स्टैंडर्ड.
  • tf32 टाइप, TensorFloat32 फ़ॉर्मैट से मेल खाता है और StableHLO में पूरी तरह से काम नहीं करता है.
• जटिल प्रकार ऐसे जटिल मान दिखाते हैं जिनका वास्तविक भाग होता है और उसी एलिमेंट टाइप का कोई काल्पनिक हिस्सा. इस्तेमाल किया जा सकने वाला कॉम्प्लेक्स टाइप 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 op, दो इनपुट वैल्यू का इस्तेमाल करता है और एक आउटपुट वैल्यू देता है. इसके मुकाबले, select_and_scatter op, तीन इनपुट वैल्यू, दो इनपुट फ़ंक्शन, और तीन इनपुट एट्रिब्यूट.

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

इनपुट फ़ंक्शन (जिन्हें अनाम फ़ंक्शन भी कहा जाता है) वे नाम वाले फ़ंक्शन से मिलते-जुलते हैं, लेकिन शामिल हैं: 1) इनमें कोई आइडेंटिफ़ायर नहीं है (इसलिए) नाम "अनाम"), 2) आउटपुट टाइप का एलान नहीं करते. फ़ंक्शन के अंदर return ऑप से अनुमान लगाया जाता है).

इनपुट फ़ंक्शन के सिंटैक्स में अभी इस्तेमाल न किया गया एक हिस्सा शामिल है (देखें Unused प्रोडक्शन, जो कि एमएलआईआर के साथ काम करने के लिए उपलब्ध है. MLIR में, "क्षेत्र" एक सामान्य कॉन्सेप्ट है, जिसे जिसमें एक से ज़्यादा "ब्लॉक" हो सकते हैं जंप ऑपरेशन के ज़रिए एक साथ जुड़े ऑपरेटर हैं. इन ब्लॉक में ऐसे आईडी हैं जो इससे जुड़े हैं Unused प्रोडक्शन में ऐसा किया गया हो, ताकि वे एक-दूसरे से अलग दिखें. StableHLO में जंप ops नहीं है, इसलिए MLIR सिंटैक्स का संबंधित हिस्सा यह है इस्तेमाल नहीं किया गया है (लेकिन अब भी मौजूद है).

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

इनपुट एट्रिब्यूट का एक नाम और एक वैल्यू होती है, जो इस्तेमाल की जा सकने वाली एट्रिब्यूट में से एक है स्थिरांक. ये प्रोग्राम के लिए स्टैटिक मेटाडेटा तय करने का मुख्य तरीका हैं एलिमेंट. उदाहरण के लिए, concatenate op, 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 ऑपरेशन, आउटपुट टाइप का अनुमान इनपुट से लगाया जा सकता है). फिर भी, ऑप हस्ताक्षर जान-बूझकर StableHLO सिंटैक्स का हिस्सा है, ताकि MLIR के साथ काम किया जा सके.

नीचे एक ऑप का उदाहरण दिया गया है जिसका स्मरणिक select_and_scatter है. यह तीन की खपत करता है इनपुट वैल्यू (%operand, %source, और %init_value), दो इनपुट फ़ंक्शन और तीन इनपुट एट्रिब्यूट (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

टेन्सर कॉन्सटेंट, नेस्ट की गई सूचियों का इस्तेमाल करके टेंसर वैल्यू दिखाते हैं 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 टेंसर पर एलिमेंट के हिसाब से abs ऑपरेशन करता है और 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]

ज़्यादा उदाहरण

जोड़ें

सिमैंटिक

दो टेंसर 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 है.
• 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] process_group में receiver.
• सभी के लिए results...@process = concatenate(operands...@process, all_gather_dim) process_group में process.

इनपुट

आउटपुट

कंस्ट्रेंट

• (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 प्रोसेस ग्रिड में हर प्रोसेस ग्रुप के अंदर, एक रिडक्शन लागू होता है हर प्रोसेस के लिए, operands टेंसर की वैल्यू के लिए computation फ़ंक्शन और 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 कहां:
  • 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.
• कॉम्प्लेक्स नंबर के लिए: कॉम्प्लेक्स 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 टेंसर बनाता है. औपचारिक तौर पर, यह ऑपरेशन को मौजूदा 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) दिए गए:
  • अगर 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) देने वाली कंपनी:
  • अगर !is_per_axis_quantized(operand) है, तो element_type(operand).
  • quantization_dimension(operand) को छोड़कर, element_type(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]
//            ]
//          ]

ज़्यादा उदाहरण

केस

सिमैंटिक

branches से सिर्फ़ एक फ़ंक्शन चलाने पर आउटपुट मिलता है यह 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]

ज़्यादा उदाहरण

सीबीआरटी

सिमैंटिक

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]

ज़्यादा उदाहरण

सील

सिमैंटिक

operand टेंसर की एलिमेंट के हिसाब से सेल करता है और result टेंसर बनाता है. आईईईई-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]

ज़्यादा उदाहरण

चोलस्की

सिमैंटिक

मैट्रिक्स के बैच के कोलेस्की अपघटन की गणना करता है.

ज़्यादा औपचारिक तौर पर, 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 को इस तरह परिभाषित किया गया है:
  • अगर 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]]

ज़्यादा उदाहरण

तुलना करें

सिमैंटिक

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 के साथ फ़्लोटिंग-पॉइंट एलिमेंट टाइप के लिए, op- मान लागू करता है आईईईई-754 के ये ऑपरेशन:

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

compare_type = TOTALORDER वाले फ़्लोटिंग-पॉइंट एलिमेंट के लिए, op इससे totalOrder और compareQuietEqual कार्रवाइयों का कॉम्बिनेशन इस्तेमाल करता है आईईईई-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 को इस तरह परिभाषित किया गया है:
  • अगर 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 वीडियो बनाया जा रहा है. कॉन्टेंट बनाने ऑपरेशन के सिमैंटिक को decomposition एट्रिब्यूट की मदद से लागू किया जाता है. कॉन्टेंट बनाने प्रोग्राम में बदलाव किए बिना, composite ऑप को इसके डिकपोज़िशन से बदला जा सकता है सिमेंटिक्स. ऐसे मामलों में जहां अपघटन को इनलाइन करने से नतीजे नहीं मिलते ऑप सिमैंटिक्स, custom_call का इस्तेमाल करना पसंद है.

version फ़ील्ड (डिफ़ॉल्ट रूप से 0 पर सेट होती है) का इस्तेमाल किसी कंपोज़िट की वैल्यू को दिखाने के लिए किया जाता है सिमेंटिक्स बदल गया है.

इनपुट

आउटपुट

कंस्ट्रेंट

• (C1) is_namespaced_op_name(name)
• (सी2) 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>

ज़्यादा उदाहरण

जोड़ें

सिमैंटिक

inputs को dimension डाइमेंशन के साथ उसी क्रम में जोड़ता है जिस क्रम में दिया गया है आर्ग्युमेंट और 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) dim(inputs..., dimension) को छोड़कर, same(shape(inputs...)).
• (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 टेंसर बनाता है.

boolean-to-any-supported-type के लिए, वैल्यू false है शून्य में बदल जाता है और true मान एक में बदल जाता है. इसके लिए any-supported-type-to-boolean रूपांतरण, शून्य मान को false और शून्य के अलावा किसी अन्य वैल्यू को true में बदल दिया जाता है. इसे देखने का तरीका नीचे दिया गया है और जटिल प्रकारों के लिए काम करते हैं.

integer-टू-इंटीजर, integer-टू-फ़्लोटिंग-पॉइंट वाले कन्वर्ज़न के लिए या 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 इस्तेमाल के दूसरे और तीसरे उदाहरण. आने वाले समय में, ये दोनों ऑपरेशन मर्ज किए जा सकते हैं 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 है, तो सभी के लिए 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) 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 और ऑपरेशन की दाईं ओर मौजूद वैल्यू में बदल दिया जाता है. प्रिसिज़न टाइप, और आउटपुट के स्टोरेज टाइप को अलग-अलग कर सकते हैं.
• accumulation_type इकट्ठा करने के लिए इस्तेमाल की गई सटीक वैल्यू.
• lhs_component_count, rhs_component_count, और num_primitive_operations तब लागू होती है, जब हम एक ऐसा एल्गोरिदम बनाते हैं जो LHS और/या RHS को कई कॉम्पोनेंट और कई सारे "प्रीमिटिव" होते हैं उन पर डॉट ऑपरेशन वैल्यू - आम तौर पर ज़्यादा सटीक जानकारी देने के लिए (उदाहरण के लिए, ज़्यादा सटीक गणना के लिए bfloat16 आर्टिफ़िशियल इंटेलिजेंस डेटाटाइप का इस्तेमाल करना: 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 के उपभोक्ता के हिसाब से एक्सेलरेटर टाइप. अगर दिया गया एल्गोरिदम, नहीं दिखाई देता है, तो कोई गड़बड़ी वैकल्पिक है. 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_non_expanding_dimensions भी स्वीकार किए जाते हैं डाइमेंशन के बढ़ते व्यवहार के बारे में अपनी अहम जानकारी ज़ाहिर करने के लिए. अगर जानकारी नहीं दी गई है, तो माना गया है कि सभी डाइमेंशन बड़े हो सकते हैं.

इनपुट

आउटपुट

कंस्ट्रेंट

• (C1) element_type(result) देने वाली कंपनी:
  • अगर !is_per_axis_quantized(operand) है, तो element_type(operand).
  • quantization_dimension(operand) को छोड़कर, element_type(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_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] दिया गया:
  • 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] दिया गया:
  • 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]]
//          ]]