ข้อกำหนดของ StableHLO

StableHLO คือชุดการดำเนินการสำหรับการดำเนินการระดับสูง (HLO) ในโมเดลแมชชีนเลิร์นนิง StableHLO ทำงานเป็นเลเยอร์ความสามารถในการพกพาระหว่างเฟรมเวิร์ก ML และคอมไพเลอร์ ML โดยเฟรมเวิร์ก ML ที่สร้างโปรแกรม StableHLO สามารถใช้งานร่วมกับคอมไพเลอร์ ML ที่ใช้โปรแกรม StableHLO

เป้าหมายของเราคือการลดความซับซ้อนและเร่งการพัฒนา ML โดยการสร้างความสามารถในการทำงานร่วมกันมากขึ้นระหว่างเฟรมเวิร์ก ML ต่างๆ (เช่น TensorFlow, JAX และ PyTorch) กับคอมไพเลอร์ ML (เช่น XLA และ IREE) ด้วยเหตุนี้ เอกสารนี้จึงระบุข้อกําหนดสําหรับภาษาโปรแกรม StableHLO

ข้อกําหนดนี้มี 3 ส่วนหลักๆ ประการแรก ส่วนโปรแกรมจะอธิบายโครงสร้างของโปรแกรม StableHLO ซึ่งประกอบด้วยฟังก์ชัน StableHLO ที่ประกอบไปด้วยการดำเนินการ StableHLO ภายในโครงสร้างดังกล่าว ส่วน Ops จะระบุความหมายของ Ops แต่ละรายการ ส่วนการดำเนินการจะระบุความหมายของการดำเนินการทั้งหมดเหล่านี้ที่ทำงานร่วมกันภายในโปรแกรม สุดท้าย ส่วนสัญลักษณ์จะอธิบายสัญลักษณ์ที่ใช้ตลอดทั้งข้อกำหนด

หากต้องการดูข้อกำหนดจาก StableHLO เวอร์ชันก่อนหน้า ให้เปิดที่เก็บรีโปที่รุ่นที่ติดแท็กที่ต้องการ เช่น ข้อกําหนดของ StableHLO v0.19.0 หากต้องการดูการเปลี่ยนแปลงที่เกิดขึ้นเมื่อมีการอัปเกรด StableHLO เป็นเวอร์ชันย่อยแต่ละเวอร์ชัน โปรดดูบันทึกเวอร์ชันใน VhloDialect.td

โปรแกรม

Program ::= {Func}

โปรแกรม StableHLO ประกอบด้วยฟังก์ชัน StableHLO กี่รายการก็ได้ ด้านล่างนี้คือตัวอย่างโปรแกรมที่มีฟังก์ชัน @main ซึ่งมีอินพุต 3 รายการ (%image, %weights และ %bias) และเอาต์พุต 1 รายการ เนื้อความของฟังก์ชันมีการดำเนินการ 6 รายการ

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

ฟังก์ชัน

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

ฟังก์ชัน StableHLO (หรือที่เรียกว่าฟังก์ชันที่มีชื่อ) ประกอบด้วยตัวระบุ อินพุต/เอาต์พุต และเนื้อหา ในอนาคต เราวางแผนที่จะเปิดตัวข้อมูลเมตาเพิ่มเติมสำหรับฟังก์ชันเพื่อให้ใช้งานร่วมกับ HLO ได้ดียิ่งขึ้น (#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} | '?'

ประเภท Tensor แสดงถึง Tensor หรืออาร์เรย์หลายมิติ โดยจะมีรูปร่างและประเภทองค์ประกอบ โดยรูปร่างจะแสดงขนาดมิติข้อมูลที่ไม่ใช่ค่าลบหรือไม่ทราบค่าตามลําดับจากน้อยไปมากของมิติข้อมูลที่เกี่ยวข้อง (หรือที่เรียกว่าแกน) ซึ่งระบุหมายเลขตั้งแต่ 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 ::= IntegerLiteral
QuantizationStorageMax ::= IntegerLiteral
QuantizationExpressedType ::= FloatType
QuantizationDimension ::= IntegerLiteral
QuantizationParameters ::= QuantizationParameter
                         | '{' QuantizationParameter {',' QuantizationParameter} '}'
QuantizationParameter ::= QuantizationScale [':' QuantizationZeroPoint]
QuantizationScale ::= FloatLiteral
QuantizationZeroPoint ::= IntegerLiteral

ประเภทองค์ประกอบที่แปลงค่าเป็นจำนวนเต็มแสดงค่าจำนวนเต็มของประเภทพื้นที่เก็บข้อมูลในช่วง storage_min ถึง storage_max (รวม) ซึ่งสอดคล้องกับค่าทศนิยมของประเภทที่แสดง สําหรับค่าจํานวนเต็ม i ที่ระบุ ระบบจะคํานวณค่าทศนิยมที่เกี่ยวข้อง f ได้เป็น f = (i - zero_point) * scale โดยที่ scale และ zero_point เรียกว่าพารามิเตอร์การแปลงค่า storage_min และ storage_max เป็นค่าที่ไม่บังคับในไวยากรณ์ แต่มีค่าเริ่มต้นเป็น min_value(storage_type) และ max_value(storage_type) ตามลำดับ องค์ประกอบประเภทที่แปลงเป็นจำนวนเต็มมีข้อจำกัดต่อไปนี้

• (C1) type(storage_min) = storage_type
• (ค2) 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 ซึ่งรวมถึงประเภท ค่า และพิจารณาว่าจะมีจุด 0 จุด 1 จุดหรืออาจเป็น 0 จุดในประเภท Tensor ที่เล็กลงหรือไม่ จากผลการพูดคุยครั้งนี้ ข้อกําหนดเกี่ยวกับจุดศูนย์อาจมีการเปลี่ยนแปลงในอนาคต (#1405)

การพูดคุยอีกเรื่องหนึ่งที่กำลังดำเนินอยู่เกี่ยวข้องกับความหมายของ QuantizationStorageMin และ QuantizationStorageMax เพื่อพิจารณาว่าควรมีข้อจำกัดสำหรับค่าเหล่านี้และค่าของ Tensor ที่แปลงค่าเป็นจำนวนเต็มหรือไม่ (#1406)

สุดท้าย เราวางแผนที่จะสำรวจการนำเสนอรูปแบบที่ไม่รู้จักและจุดเริ่มต้น เช่นเดียวกับที่เราวางแผนที่จะสำรวจการนำเสนอขนาดมิติข้อมูลที่ไม่ทราบ (#1407)

ประเภท tensor ที่กำหนดแสดงถึง Tensor ที่มีองค์ประกอบที่เล็กลง tensor เหล่านี้เหมือนกับ tensor ปกติทุกประการ เว้นแต่องค์ประกอบของ Tensor จะมีประเภทองค์ประกอบที่เล็กลงแทนที่จะเป็นประเภทองค์ประกอบปกติ

ในเทนเซอร์ที่ผ่านการแปลงค่า การแปลงค่าอาจเป็นต่อเทนเซอร์ ซึ่งหมายความว่ามี scale และ zero_point 1 คู่สำหรับเทนเซอร์ทั้งรายการ หรืออาจเป็นต่อแกน ซึ่งหมายความว่ามี scales และ zero_points หลายคู่ โดย 1 คู่ต่อส่วนย่อยของมิติข้อมูล quantization_dimension ที่เฉพาะเจาะจง กล่าวอย่างเป็นทางการคือ ในเทนเซอร์ t ที่มีการปรับขนาดตามแกน จะมี dim(t, quantization_dimension) ส่วนของ quantization_dimension ดังนี้ t[:, ..., 0, ..., :], t[:, ..., 1, ..., :] เป็นต้น องค์ประกอบทั้งหมดในส่วนที่ i ใช้ scales[i] และ zero_points[i] เป็นพารามิเตอร์การปรับขนาด ประเภท Tensor ที่แปลงค่าเป็นจำนวนเต็มมีข้อจำกัดต่อไปนี้

• สำหรับการปันส่วนต่อ Tensor ให้ทำดังนี้
  • ไม่มีข้อจำกัดเพิ่มเติม
• สําหรับการแจกแจงต่อแกน ให้ทำดังนี้
  • (C12) quantization_dimension < rank(self)
  • (C13) dim(self, quantization_dimension) = size(scales)

TokenType ::= 'token'

ประเภทโทเค็นแสดงถึงโทเค็น ซึ่งเป็นค่าที่ไม่โปร่งใสที่การดำเนินการบางอย่างสร้างขึ้นและใช้ โทเค็นใช้เพื่อกำหนดลําดับการดําเนินการของการดำเนินการตามที่อธิบายไว้ในส่วนการดําเนินการ

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

ประเภททูเปิลแสดงถึงทูเปิล เช่น รายการที่หลากหลาย Tuples เป็นฟีเจอร์เดิมที่มีไว้เพื่อใช้งานร่วมกับ HLO เท่านั้น ใน HLO ระบบจะใช้ทูเพลเพื่อแสดงอินพุตและเอาต์พุตแบบผันแปร ใน StableHLO ระบบจะรองรับอินพุตและเอาต์พุตแบบผันแปรโดยค่าเริ่มต้น และการใช้ทูเพลตใน StableHLO มีไว้เพื่อแสดง HLO ABI อย่างครอบคลุมเท่านั้น เช่น T, tuple<T> และ tuple<tuple<T>> อาจแตกต่างกันอย่างมากโดยขึ้นอยู่กับการใช้งาน ในอนาคต เราวางแผนที่จะเปลี่ยนแปลง HLO ABI ซึ่งอาจทำให้เรานำประเภททูเปิลออกจาก StableHLO ได้ (#598)

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

ประเภทองค์ประกอบแสดงองค์ประกอบของประเภท Tensor ภาษาประเภทนี้ต่างจากภาษาโปรแกรมหลายภาษา ตรงไม่ใช่คลาสแรกใน 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 โดยรวม
• ประเภทตัวเลขทศนิยมอาจเป็นอย่างใดอย่างหนึ่งต่อไปนี้
  • ตัวเลขทศนิยม 8 บิต f8E3M4, f8E4M3 และ f8E5M2 ตามรูปแบบ IEEE-754
  • ประเภท f8E4M3FN และ f8E5M2 สอดคล้องกับการเข้ารหัส E4M3 และ E5M2 ของรูปแบบ FP8 ที่อธิบายไว้ในรูปแบบ FP8 สําหรับการเรียนรู้เชิงลึก
  • ประเภท f8E4M3FNUZ และ f8E5M2FNUZ ที่สอดคล้องกับการเข้ารหัส E4M3 และ E5M2 ของรูปแบบ FP8 ที่อธิบายไว้ในรูปแบบตัวเลข 8 บิตสําหรับเครือข่ายประสาทเทียมลึก
  • f8E4M3B11FNUZ ประเภทที่สอดคล้องกับการเข้ารหัส E4M3 ของรูปแบบ FP8 ที่อธิบายไว้ในการฝึกและการอนุมาน 8 บิตแบบไฮบริด (HFP8) สำหรับเครือข่ายระบบประสาทแบบลึก
  • ประเภท bf16 ที่สอดคล้องกับรูปแบบ bfloat16 ที่อธิบายไว้ใน BFloat16: เคล็ดลับสู่ประสิทธิภาพสูงบน Cloud TPU
  • ประเภท f16, f32 และ f64 สอดคล้องกับรูปแบบ binary16 ("ครึ่งความแม่นยำ"), binary32 ("ความแม่นยำเดี่ยว") และ binary64 ("ความแม่นยำคู่") ตามลำดับที่อธิบายไว้ในมาตรฐาน IEEE 754
  • ประเภท tf32 สอดคล้องกับรูปแบบ TensorFloat32 และรองรับใน StableHLO แบบจํากัด
  • f4E2M1FN, f6E2M3FN, f6E3M2FN และ f8E8M0FNU MX (Microscaling) ประเภท ที่อธิบายไว้ใน ข้อมูลจำเพาะของรูปแบบ Microscaling ของ OCP
• ประเภทที่ซับซ้อนหมายถึงค่าเชิงซ้อนที่มีส่วนจริงและส่วนจินตภาพขององค์ประกอบประเภทเดียวกัน ประเภทเชิงซ้อนที่รองรับคือ complex<f32> (ทั้ง 2 ส่วนเป็นประเภท f32) และ complex<f64> (ทั้ง 2 ส่วนเป็นประเภท f64)

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

ประเภทฟังก์ชันแสดงทั้งฟังก์ชันที่มีชื่อและฟังก์ชันที่ไม่มีชื่อ ฟังก์ชันมีประเภทอินพุต (รายการประเภททางด้านซ้ายของ ->) และประเภทเอาต์พุต (รายการประเภททางด้านขวาของ ->) ในภาษาโปรแกรมหลายภาษา ประเภทฟังก์ชันจะเป็นประเภทระดับบนสุด แต่ไม่ใช่ใน StableHLO

StringType ::= 'string'

ประเภทสตริงแสดงลำดับไบต์ ประเภทสตริงไม่ใช่คลาสแรกใน StableHLO และจะใช้เพื่อระบุข้อมูลเมตาแบบคงที่สำหรับองค์ประกอบของโปรแกรมเท่านั้น ซึ่งต่างจากในภาษาโปรแกรมหลายภาษา

การดำเนินการ

การดำเนินการ StableHLO (หรือที่เรียกว่าการดำเนินการ) หมายถึงชุดการดำเนินการระดับสูงแบบปิดในโมเดลแมชชีนเลิร์นนิง ตามที่เราได้กล่าวไปข้างต้น ไวยากรณ์ของ StableHLO ได้รับแรงบันดาลใจอย่างมากจาก MLIR ซึ่งไม่ใช่ทางเลือกที่เหมาะกับการใช้งานมากที่สุด แต่อาจกล่าวได้ว่าเหมาะกับเป้าหมายของ StableHLO ในการสร้างความสามารถในการทำงานร่วมกันมากขึ้นระหว่างเฟรมเวิร์ก ML กับคอมไพเลอร์ ML

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 ใช้ค่าอินพุต 2 ค่าและสร้างค่าเอาต์พุต 1 ค่า ในทางกลับกัน select_and_scatter op จะใช้ค่าอินพุต 3 ค่า ฟังก์ชันอินพุต 2 รายการ และแอตทริบิวต์อินพุต 3 รายการ

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

ฟังก์ชันอินพุต (หรือที่เรียกว่าฟังก์ชันนิรนาม) คล้ายกับฟังก์ชันที่มีชื่อมาก ยกเว้นว่า 1) ไม่มีตัวระบุ (จึงเป็นที่มาของชื่อ "นิรนาม") 2) ไม่ประกาศประเภทเอาต์พุต (ระบบจะอนุมานประเภทเอาต์พุตจาก return op ภายในฟังก์ชัน)

รูปแบบคำสั่งสำหรับฟังก์ชันอินพุตมีส่วนที่ไม่ได้ใช้งานในปัจจุบัน (ดูUnusedเวอร์ชันที่ใช้งานจริงด้านบน) ซึ่งมีไว้เพื่อใช้งานร่วมกับ MLIR ใน MLIR จะมีแนวคิด "ภูมิภาค" ทั่วไปมากขึ้น ซึ่งอาจมี "บล็อก" ของการดำเนินการหลายรายการที่เชื่อมต่อกันผ่านการดำเนินการ Jump บล็อกเหล่านี้มีรหัสที่สอดคล้องกับUnusedเวอร์ชันที่ใช้งานจริงเพื่อให้แยกความแตกต่างกันได้ StableHLO ไม่มีการข้าม ดังนั้นระบบจึงไม่ได้ใช้ส่วนที่เกี่ยวข้องของไวยากรณ์ MLIR (แต่ยังคงอยู่)

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

แอตทริบิวต์อินพุตมีชื่อและค่าซึ่งเป็นหนึ่งในค่าคงที่ที่รองรับ ซึ่งเป็นวิธีหลักในการระบุข้อมูลเมตาแบบคงที่สำหรับองค์ประกอบโปรแกรม เช่น อ็อป concatenate ใช้แอตทริบิวต์ dimension เพื่อระบุมิติข้อมูลที่จะใช้ต่อท้ายค่าอินพุต ในทํานองเดียวกัน slice op ใช้แอตทริบิวต์หลายรายการ เช่น 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 ส่วนใหญ่ ประเภทเอาต์พุตจะอนุมานได้จากอินพุต) อย่างไรก็ตาม ลายเซ็น op ก็ตั้งใจเป็นส่วนหนึ่งของไวยากรณ์ StableHLO เพื่อความเข้ากันได้กับ MLIR

ด้านล่างนี้คือตัวอย่างการดำเนินการที่มีคําช่วยจําคือ select_and_scatter ระบบจะใช้ค่าอินพุต 3 ค่า (%operand, %source และ %init_value), ฟังก์ชันอินพุต 2 รายการ และแอตทริบิวต์อินพุต 3 รายการ (window_dimensions, window_strides และ padding) โปรดทราบว่าลายเซ็นของการดำเนินการมีเฉพาะประเภทของค่าอินพุต (แต่ไม่ใช่ประเภทของฟังก์ชันอินพุตและแอตทริบิวต์ที่ป้อนในบรรทัด)

%result = "stablehlo.select_and_scatter"(%operand, %source, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "stablehlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GE>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "stablehlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "stablehlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "stablehlo.return"(%0) : (tensor<i32>) -> ()
}) {
  window_dimensions = dense<[3, 1]> : tensor<2xi64>,
  window_strides = dense<[2, 1]> : tensor<2xi64>,
  padding = dense<[[0, 1], [0, 0]]> : tensor<2x2xi64>
} : (tensor<4x2xi32>, tensor<2x2xi32>, tensor<i32>) -> tensor<4x2xi32>

ค่าคงที่

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

ค่าคงที่ StableHLO มีนิพจน์และประเภทที่ร่วมกันแสดงค่า StableHLO โดยทั่วไป ประเภทเป็นส่วนหนึ่งของไวยากรณ์ของค่าคงที่ ยกเว้นในกรณีที่ไม่คลุมเครือ (เช่น ค่าคงที่บูลีนมีประเภท i1 อย่างชัดเจน ในขณะที่ค่าคงที่จำนวนเต็มอาจมีประเภทได้หลายประเภท)

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

ค่าคงที่บูลีนแสดงค่าบูลีน true และ false ค่าคงที่บูลีนมีประเภท i1

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

ค่าคงที่จำนวนเต็มแสดงค่าจำนวนเต็มผ่านสตริงที่ใช้การเขียนเลขทศนิยมหรือเลขฐาน 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]

ค่าคงที่ทศนิยมแสดงค่าทศนิยมผ่านสตริงที่ใช้สัญกรณ์ทศนิยมหรือสัญกรณ์วิทยาศาสตร์ นอกจากนี้ คุณยังใช้การเขียนเลขฐานสิบหกเพื่อระบุบิตพื้นฐานในรูปแบบทศนิยมของประเภทที่เกี่ยวข้องได้โดยตรง ค่าคงที่ทศนิยมมีข้อจำกัดต่อไปนี้

• (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))
• (ค2) is_wellformed(imaginary_part, complex_element_type(complex_type))

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

ค่าคงที่ของ Tensor แสดงค่า Tensor โดยใช้ลิสต์ที่ซ้อนกันซึ่งระบุผ่านรูปแบบนิพจน์ NumPy เช่น dense<[[1, 2, 3], [4, 5, 6]]> : tensor<2x3xi32> represent a tensor value with the following mapping from indices to elements: {0, 0} => 1, {0, 1} => 2, {0, 2} => 3, {1, 0} => 4, {1, 1} => 5, {1, 2} => 6 ลำดับการจัดเก็บองค์ประกอบเหล่านี้ในหน่วยความจำจะกำหนดการใช้งาน ค่าคงที่ Tensor มีข้อจำกัดต่อไปนี้

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

ความหมาย

ดำเนินการ abs ตามองค์ประกอบในเทนเซอร์ operand และสร้างเทนเซอร์ result ทําสิ่งต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สำหรับจำนวนเต็มที่มีเครื่องหมาย: โมดูลัสจำนวนเต็ม
• สำหรับตัวเลขทศนิยม: abs จาก IEEE-754
• สำหรับจำนวนเชิงซ้อน: โมดูลัสเชิงซ้อน
• สำหรับประเภทที่มีการแปลงเป็นจำนวนเต็ม: dequantize_op_quantize(abs, operand, type(result))

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) shape(result) = shape(operand)
• (C2) baseline_element_type(result) มีคำจำกัดความดังนี้
  • complex_element_type(element_type(operand)) หาก is_complex(operand)
  • baseline_element_type(operand) ในกรณีอื่น

ตัวอย่าง

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

 ตัวอย่างเพิ่มเติม

เพิ่ม

ความหมาย

ทำการบวก Tensor 2 ตัว lhs และ rhs และสร้าง Tensor result ทําสิ่งต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สําหรับบูลีน: OR เชิงตรรกะ
• สําหรับจํานวนเต็ม: การเพิ่มจํานวนเต็ม
• สำหรับตัวเลขทศนิยม: addition จาก IEEE-754
• สำหรับจำนวนเชิงซ้อน: การบวกเชิงซ้อน
• สำหรับประเภทที่มีการแปลงเป็นจำนวนเต็ม: dequantize_op_quantize(add, lhs, rhs, type(result))

อินพุต

เอาต์พุต

ข้อจำกัด

• หากการดำเนินการใช้เทนเซอร์ที่ไม่ผ่านการแปลงค่า ให้ทำดังนี้
  • (C1) type(lhs) = type(rhs) = type(result)
• หากการดำเนินการใช้ Tensor ที่เล็กลง
  • (ค2) 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 ให้ต่อค่าของเทอร์เซอร์ operands จากแต่ละกระบวนการตาม all_gather_dim และสร้างเทอร์เซอร์ results

การดำเนินการนี้จะแบ่งตารางกริดของกระบวนการ StableHLO เป็น process_groups ซึ่งได้กำหนดไว้ดังต่อไปนี้

• cross_replica(replica_groups) if 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) if 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...)
• (ค2) is_unique(replica_groups)
• (C3) size(replica_groups) มีคำจำกัดความดังนี้
  • num_replicas หากใช้ cross_replica
  • num_replicas หากใช้ cross_replica_and_partition
  • num_processes หากใช้ flattened_ids
• (C4) 0 <= replica_groups < size(replica_groups)
• (C5) หากเป็น use_global_device_ids = true ให้ channel_id > 0
• (C6) type(results...) = type(operands...) ยกเว้นในกรณีต่อไปนี้
  • dim(results..., all_gather_dim) = dim(operands..., all_gather_dim) * dim(process_groups, 1)

ตัวอย่าง

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

 ตัวอย่างเพิ่มเติม

all_reduce

ความหมาย

ภายในกลุ่มกระบวนการแต่ละกลุ่มในตารางกริดกระบวนการ StableHLO ให้ใช้ฟังก์ชันการลด computation กับค่าของเทนเซอร์ operands จากแต่ละกระบวนการ และสร้างเทนเซอร์ results

การดำเนินการนี้จะแบ่งตารางกริดกระบวนการ StableHLO เป็น process_groups ซึ่งกําหนดไว้ดังนี้

• cross_replica(replica_groups) if 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) if 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) มีความหมายดังต่อไปนี้
  • num_replicas หากใช้ cross_replica
  • num_replicas หากใช้ cross_replica_and_partition
  • num_processes หากใช้ flattened_ids
• (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...)
• (ค7) 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 ซึ่งได้กำหนดไว้ดังต่อไปนี้

• cross_replica(replica_groups) หาก channel_id <= 0
• cross_partition(replica_groups) หาก channel_id > 0

หลังจากนั้น ให้ทำดังนี้ภายในแต่ละ process_group

• split_parts...@sender = split(operands...@sender, split_count, split_dimension) สำหรับ sender ทั้งหมดใน process_group
• 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) มีความหมายดังนี้
  • num_replicas หากใช้ cross_replica
  • num_partitions หากใช้ cross_partition
• (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]]

 ตัวอย่างเพิ่มเติม

และ

ความหมาย

ดำเนินการ AND ตามองค์ประกอบที่มี 2 Tensor คือ lhs และ rhs และสร้าง Tensor result ดำเนินการต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สําหรับบูลีน: AND เชิงตรรกะ
• สำหรับจำนวนเต็ม: Bitwise 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

ความหมาย

ดำเนินการ atan2 ตามองค์ประกอบในเทนเซอร์ lhs และ rhs และสร้างเทนเซอร์ result ดำเนินการต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สำหรับแบบลอย: atan2 จาก IEEE-754
• สำหรับจำนวนเชิงซ้อน: complex atan2
• สำหรับประเภทที่เล็กลง: dequantize_op_quantize(atan2, lhs, rhs, type(result))

อินพุต

เอาต์พุต

ข้อจำกัด

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

ตัวอย่าง

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

ตัวอย่างเพิ่มเติม

batch_norm_grad

ความหมาย

ประมวลผลการไล่ระดับสีของอินพุตหลายรายการของ batch_norm_training Backpropagating จาก grad_output และสร้าง grad_operand, grad_scale และ grad_offset tensor การดำเนินการนี้สามารถแสดงเป็นการจัดเรียงการดำเนินการ 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

ความหมาย

ปรับค่า Tensor operand ให้เป็นมาตรฐานในทุกมิติข้อมูลยกเว้นมิติข้อมูล feature_index และสร้าง Tensor 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)
• (ค7) 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 รูปแบบที่เป็นทางการมากขึ้น อาจแสดงเป็นอนุพันธ์ของการดำเนินการ StableHLO ที่มีอยู่โดยใช้ไวยากรณ์ Python ดังนี้

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

อรรถศาสตร์

ดำเนินการบิตแคสต์บน tensor ของ operand และสร้าง tensor ของ result ซึ่งมีการตีความบิตของ Tensor ทั้ง operand อีกครั้งโดยใช้ประเภทของ tensor ของ 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 จะแสดงผลการแสดงค่าหนึ่งๆ ในหน่วยความจำ และระบบจะกำหนดลักษณะการทำงานของค่าดังกล่าวเนื่องจากมีการกำหนดการแสดง Tensor ที่แน่นอนและมีการกำหนดการใช้งานการแสดงประเภทองค์ประกอบที่แน่นอนด้วย

อินพุต

เอาต์พุต

ข้อจำกัด

• (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
  • dim(result, i) = dim(operand, i) สำหรับ 0 <= i < R ทั้งหมด
  • dim(result, R) * num_bits(E') = num_bits(E)
  • หาก num_bits(E') > num_bits(E):
  • rank(result) = R - 1
  • dim(result, i) = dim(operand, i) สำหรับ 0 <= i < R ทั้งหมด
  • dim(operand, R - 1) * num_bits(E) = num_bits(E')
• (C2) หากเป็น is_complex(operand) or is_complex(result) แล้ว is_complex(operand) and is_complex(result)

ตัวอย่าง

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

 ตัวอย่างเพิ่มเติม

broadcast_in_dim

ความหมาย

ขยายมิติข้อมูลและ/หรืออันดับของเทมพอร์ลอินพุตโดยการทำซ้ำข้อมูลในเทมพอร์ล operand และสร้างเทมพอร์ล result เขียนเป็นทางการคือ result[result_index] = operand[operand_index] โดยที่สำหรับ d ทั้งหมดใน axes(operand)

• operand_index[d] = 0 หาก dim(operand, d) = 1
• จ่าย operand_index[d] = result_index[broadcast_dimensions[d]]

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) element_type(result) คำนวณจากข้อมูลต่อไปนี้
  • element_type(operand) หาก !is_per_axis_quantized(operand)
  • element_type(operand) ยกเว้น quantization_dimension(operand), scales(operand) และ zero_points(operand) อาจแตกต่างจาก quantization_dimension(result), scales(result) และ zero_points(result) ตามลำดับ
• (ค2) size(broadcast_dimensions) = rank(operand)
• (C3) 0 <= broadcast_dimensions < rank(result)
• (C4) is_unique(broadcast_dimensions)
• (C5) สําหรับ d ทั้งหมดใน axes(operand) ให้ทําดังนี้
  • 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 เพียง 1 รายการเท่านั้น โดยขึ้นอยู่กับค่าของ index เขียนเป็นทางการคือ result = selected_branch() โดยที่

• selected_branch = branches[index] หากเป็น 0 <= index < size(branches)
• จ่าย 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]

 ตัวอย่างเพิ่มเติม

CBT

ความหมาย

ดำเนินการรากคิวบิกทีละองค์ประกอบในเทนเซอร์ operand และสร้างเทนเซอร์ result ทําสิ่งต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สำหรับตัวเลขทศนิยม: rootn(x, 3) จาก IEEE-754
• สำหรับจำนวนเชิงซ้อน: รากที่สามเชิงซ้อน
• สำหรับประเภทที่เล็กลง: dequantize_op_quantize(cbrt, operand, type(result))

อินพุต

เอาต์พุต

ข้อจำกัด

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

ตัวอย่าง

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

 ตัวอย่างเพิ่มเติม

ceil

ความหมาย

ดำเนินการปัดเศษทศนิยมขึ้นทีละองค์ประกอบของเทนเซอร์ operand และสร้างเทนเซอร์ result ใช้การดำเนินการ roundToIntegralTowardPositive จากข้อกำหนด IEEE-754 สําหรับประเภทที่กําหนดค่าได้ ให้ทําการ 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 ของเมทริกซ์หลายรายการ

พูดให้ถูกต้องมากขึ้นคือ สำหรับ i ทั้งหมดใน index_space(result) result[i0, ..., iR-3, :, :] เป็นการแยก Cholesky ของ a[i0, ..., iR-3, :, :] ในรูปแบบของเมทริกซ์สามเหลี่ยมล่าง (หาก lower เป็น true) หรือสามเหลี่ยมบน (หาก lower เป็น false) ค่าเอาต์พุตในสามเหลี่ยมตรงข้าม ซึ่งก็คือสามเหลี่ยมบนแบบเข้มงวดหรือสามเหลี่ยมล่างแบบเข้มงวดตามลำดับนั้นกำหนดโดยการใช้งาน

หากมี i ที่เมทริกซ์อินพุตไม่ใช่เมทริกซ์ Hermitian ที่แน่นอนเชิงบวก ระบบจะไม่ระบุลักษณะการทำงาน

สําหรับประเภทที่กําหนดค่าได้ ให้ทําการ dequantize_op_quantize(lambda operand: cholesky(operand, lower), a, type(result))

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) baseline_type(a) = baseline_type(result)
• (ค2) 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]
//          ]

หนีบ

อรรถศาสตร์

ยึดทุกองค์ประกอบของ tensor ของ operand ระหว่างค่าต่ำสุดและสูงสุด และสร้าง Tensor ขนาด 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)
• (ค2) 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 ซึ่งกําหนดไว้ดังนี้

• cross_replica(replica_groups) หากเป็น channel_id <= 0
• cross_partition(replica_groups) หาก channel_id > 0

หลังจากนั้น 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 หมายถึง
  • num_replicas หากใช้ cross_replica
  • num_partitions หากใช้ cross_partition
• (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 ให้ส่งค่าของ Tensor operand จากกระบวนการต้นทางไปยังกระบวนการเป้าหมายและสร้าง Tensor result

การดำเนินการนี้จะแบ่งตารางกริดกระบวนการ StableHLO เป็น process_groups ซึ่งกําหนดไว้ดังนี้

• cross_replica(source_target_pairs) หาก channel_id <= 0
• cross_partition(source_target_pairs) หาก channel_id > 0

หลังจากนั้น 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
• (ค2) is_unique(source_target_pairs[:, 0])
• (C3) is_unique(source_target_pairs[:, 1])
• (C4) 0 <= source_target_pairs < N โดยที่ N มีความหมายดังนี้
  • num_replicas หากใช้ cross_replica
  • num_partitions หากใช้ cross_partition
• (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]]

 ตัวอย่างเพิ่มเติม

เปรียบเทียบ

ความหมาย

ทำการเปรียบเทียบ tensor ของ lhs และ rhs ตามองค์ประกอบ comparison_direction และ compare_type และสร้าง Tensor เป็น 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)
• (ค2) shape(lhs) = shape(rhs) = shape(result)
• (C3) compare_type มีคำจำกัดความดังนี้
  • SIGNED หาก is_signed_integer(element_type(lhs))
  • UNSIGNED หากเป็น is_unsigned_integer(element_type(lhs)) or is_boolean(element_type(lhs))
  • FLOAT หรือ TOTALORDER หาก is_float(element_type(lhs))
  • FLOAT หาก is_complex(element_type(lhs))

ตัวอย่าง

// %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)
• (ค2) 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)
• (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

ความหมาย

ต่อ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) same(shape(inputs...)) ยกเว้น dim(inputs..., dimension)
• (C3) 0 < size(inputs)
• (C4) 0 <= dimension < rank(inputs[0])
• (C5) element_type(result) = element_type(inputs[0])
• (C6) shape(result) = shape(inputs[0]) ยกเว้นในกรณีต่อไปนี้
  • dim(result, dimension) = dim(inputs[0], dimension) + ...

ตัวอย่าง

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

 ตัวอย่างเพิ่มเติม

ค่าคงที่

ความหมาย

สร้างเทนเซอร์ output จากค่าคงที่ value

อินพุต

เอาต์พุต

ข้อจำกัด

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

 ตัวอย่างเพิ่มเติม

ทำให้เกิด Conversion

ความหมาย

ทำการแปลงทีละองค์ประกอบจากองค์ประกอบประเภทหนึ่งไปยังอีกประเภทหนึ่งในเทนเซอร์ operand และสร้างเทนเซอร์ result

สําหรับการเปลี่ยนค่า boolean เป็นประเภทที่รองรับ ระบบจะแปลงค่า false เป็น 0 และแปลงค่า true เป็น 1 สําหรับการแปลง any-supported-type-to-boolean ระบบจะแปลงค่า 0 เป็น false และแปลงค่าที่ไม่ใช่ 0 เป็น true ดูวิธีการทํางานของการดำเนินการนี้กับประเภทที่ซับซ้อนได้ที่ด้านล่าง

สำหรับ Conversion ที่เกี่ยวข้องกับ integer-to-integer, integer-to-Floing-point หรือ Floating-point-to-Floing-point หากค่าต้นทางสามารถแสดงในประเภทปลายทางได้อย่างถูกต้อง ค่าผลลัพธ์จะเป็นค่าการเป็นตัวแทนที่ถูกต้อง มิฉะนั้น ระบบจะกำหนดลักษณะการทำงานในภายหลัง (#180)

สําหรับการแปลงที่เกี่ยวข้องกับทศนิยมเป็นจํานวนเต็ม ระบบจะตัดเศษส่วน หากค่าที่ถูกตัดให้สั้นลงแสดงในประเภทปลายทางไม่ได้ ลักษณะการทํางานจะยังไม่ทราบ (#180)

Conversion ที่เกี่ยวข้องกับ Conversion แบบซับซ้อนไปซับซ้อนมีลักษณะการทำงานเดียวกันกับ Conversion จุดลอยตัวไปยังจุดลอยตัว สำหรับการแปลงส่วนจริงและส่วนจินตภาพ

สําหรับการเปลี่ยนประเภทจากเชิงซ้อนเป็นประเภทอื่นๆ และประเภทอื่นๆ เป็นเชิงซ้อน ระบบจะไม่สนใจค่าจินตภาพของต้นทางหรือตั้งค่าจินตภาพของปลายทางเป็น 0 ตามลําดับ การแปลงของส่วนที่เป็นจำนวนจริงจะเป็นไปตามการแปลงทศนิยม

โดยหลักการแล้ว การดำเนินการนี้อาจแสดงการลดจำนวน (Conversion จาก Tensor ที่เล็กลงเป็น Tensor ปกติ) การแปลงปริมาณ (การแปลงจาก Tensor ปกติเป็น Tensor ที่ควอนไซส์) และการแปลงจำนวนใหม่ (การแปลงระหว่าง Tensor ที่เล็กลง) แต่ในตอนนี้เรามีการดำเนินการเฉพาะสำหรับ -uniform_dequantize สำหรับ Use Case แรกและ 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] where 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)
• (ค2) 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
• (ค7) 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) มีคำจำกัดความดังนี้
  • dim(lhs, input_batch_dimension) / batch_group_count หากเป็น result_dim = output_batch_dimension
  • dim(rhs, kernel_output_feature_dimension) หาก result_dim = 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]]
//          ]]

 ตัวอย่างเพิ่มเติม

โคไซน์

ความหมาย

ดำเนินการโคไซน์ตามองค์ประกอบบน tensor ของ operand และสร้าง Tensor result ทําสิ่งต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สำหรับตัวเลขทศนิยม: cos จาก IEEE-754
• สำหรับจำนวนเชิงซ้อน: โคไซน์เชิงซ้อน
• สำหรับประเภทที่มีการแปลงเป็นจำนวนเต็ม: 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

ความหมาย

นับจํานวนบิต 0 นําหน้าขององค์ประกอบในเทนเซอร์ 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 ทําสิ่งต่อไปนี้โดยขึ้นอยู่กับประเภทองค์ประกอบ

• สําหรับจํานวนเต็ม: การหารจํานวนเต็มซึ่งให้ผลหารพีชคณิตโดยทิ้งเศษส่วน
• สำหรับแบบลอย: division จาก IEEE-754
• สำหรับจำนวนเชิงซ้อน: การหารจำนวนเชิงซ้อน
• สำหรับประเภทที่มีการแปลงค่าเป็นจำนวนเต็ม ให้ทำดังนี้
  • 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 และสร้าง Tensor 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 จะควบคุมการแลกเปลี่ยนระหว่างความเร็วและความแม่นยำสำหรับการคำนวณในแบ็กเอนด์ Accelerator ค่านี้อาจเป็นค่าใดค่าหนึ่งต่อไปนี้ (ขณะนี้ ความหมายของค่า enum เหล่านี้ยังไม่ได้รับการระบุอย่างละเอียด แต่เราวางแผนที่จะแก้ไขปัญหานี้ใน #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 ออกเป็นคอมโพเนนต์หลายรายการ และดำเนินการ DOT "พื้นฐาน" หลายรายการกับค่าเหล่านั้น ซึ่งมักจะเป็นการจําลองความแม่นยําที่สูงขึ้น (เช่น การใช้ประโยชน์จากประเภทข้อมูลปัญญาประดิษฐ์ 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)
• (ค2) 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)
• (ค7) 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_dimop แต่จะมีการระบุรูปแบบผลลัพธ์แบบไดนามิกผ่าน output_dimensions

นอกจากนี้ การดำเนินการยังยอมรับแอตทริบิวต์ที่ไม่บังคับ known_expanding_dimensions, known_nonexpanding_dimensions เพื่อแสดงความรู้แบบคงที่เกี่ยวกับลักษณะการขยายตัวของมิติข้อมูล หากไม่ได้ระบุ ระบบจะถือว่ามิติข้อมูลทั้งหมดอาจขยาย

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) element_type(result) คำนวณจากข้อมูลต่อไปนี้
  • element_type(operand) หาก !is_per_axis_quantized(operand)
  • element_type(operand) ยกเว้น quantization_dimension(operand), scales(operand) และ zero_points(operand) อาจแตกต่างจาก quantization_dimension(result), scales(result) และ zero_points(result) ตามลำดับ
• (ค2) size(broadcast_dimensions) = rank(operand)
• (C3) 0 <= broadcast_dimensions < rank(result)
• (C4) is_unique(broadcast_dimensions)
• (C5) สําหรับ d ทั้งหมดใน axes(operand) ให้ทําดังนี้
  • dim(operand, d) = 1หรือ
  • dim(operand, d) = dim(result, broadcast_dimensions[d])
• (C6) หาก is_per_axis_quantized(result)
  • quantization_dimension(result) = broadcast_dimensions[quantization_dimension(operand)]
  • ถ้าdim(operand, quantization_dimension(operand)) = 1 ให้ scales(result)[i] = scales(operand)[0] and zero_points(result)[i] = zero_points(operand)[0] for i in range(dim(result, quantization_dimension(result)))
• (C7) size(output_dimensions) = rank(result)
• (C8) is_unique(known_expanding_dimensions + known_nonexpanding_dimensions)
• (C9) 0 <= known_expanding_dimensions < rank(operand)
• (C10) 0 <= known_nonexpanding_dimensions < rank(operand)

ตัวอย่าง

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

 ตัวอย่างเพิ่มเติม

dynamic_conv

ความหมาย

การดำเนินการนี้มีฟังก์ชันการทำงานเหมือนกับ convolution op แต่จะมีการระบุการเติมข้อมูลแบบไดนามิกผ่าน padding

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) N = rank(lhs) = rank(rhs)
• (ค2) 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
• (ค7) 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) มีคำจำกัดความดังนี้
  • dim(lhs, input_batch_dimension) / batch_group_count หากเป็น result_dim = output_batch_dimension
  • dim(rhs, kernel_output_feature_dimension) หาก result_dim = 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]]
//          ]]