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

StableHLO คือชุดการดำเนินการสำหรับการดำเนินการระดับสูง (HLO) ในโมเดลแมชชีนเลิร์นนิง (ML) 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 | BufferType
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 มิติ (หรืออีกนัยหนึ่งคือ 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
• (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 รวมถึงประเภท ค่า และความเป็นไปได้ที่จะมีจุดศูนย์เพียงจุดเดียวหรือ อาจมีหลายจุดในประเภทเทนเซอร์ที่ควอนไทซ์ จาก ผลการอภิปรายนี้ ข้อกำหนดเกี่ยวกับคะแนน 0 อาจมีการเปลี่ยนแปลง ในอนาคต (#1405)

การอภิปรายอีกเรื่องที่กำลังดำเนินอยู่เกี่ยวข้องกับความหมายของ QuantizationStorageMin และ QuantizationStorageMax เพื่อพิจารณาว่าควรมีการกำหนดข้อจำกัดใดๆ กับค่าเหล่านี้และค่าของเทนเซอร์ที่ผ่านการหาปริมาณหรือไม่ (#1406)

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

ประเภทเทนเซอร์ที่เล็กลงแสดงเทนเซอร์ที่มีองค์ประกอบที่เล็กลง เทนเซอร์เหล่านี้เหมือนกับเทนเซอร์ปกติทุกประการ ยกเว้นว่าองค์ประกอบของเทนเซอร์มีประเภทองค์ประกอบที่ผ่านการวัดปริมาณแทนที่จะเป็นประเภทองค์ประกอบปกติ

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

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

TokenType ::= 'token'

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

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

ประเภทบัฟเฟอร์แสดงถึงบัฟเฟอร์ ตัวอย่างเช่น ใน XLA บัฟเฟอร์คือ อาร์เรย์หลายมิติที่มีพื้นที่เก็บข้อมูลที่สอดคล้องกัน เช่นเดียวกับประเภทเทนเซอร์ ประเภทบัฟเฟอร์มีรูปร่างและประเภทองค์ประกอบ โดยรูปร่างแสดงถึง ขนาดมิติข้อมูลที่ไม่เป็นลบหรือไม่รู้จักตามลำดับจากน้อยไปมากของ มิติข้อมูลที่เกี่ยวข้อง (เรียกอีกอย่างว่าแกน) ซึ่งมีหมายเลขตั้งแต่ 0 ถึง R-1 จำนวนมิติข้อมูล R เรียกว่าอันดับ เช่น memref<2x3xf32> เป็นประเภทบัฟเฟอร์ที่มีรูปร่าง 2x3 และประเภทองค์ประกอบ f32 มี 2 มิติ (หรืออีกนัยหนึ่งคือ 2 แกน) ได้แก่ มิติที่ 0 และมิติที่ 1 ซึ่งมีขนาดเป็น 2 และ 3 โดยมีอันดับที่ 2

คุณจัดสรรบัฟเฟอร์ได้โดยใช้ custom_call ถึง CreateBuffer หรือ Pin และยกเลิกการจัดสรรผ่าน custom_call ถึง Unpin มีเพียงcustom_call ops เท่านั้นที่อ่านและ เขียนเนื้อหาภายในบัฟเฟอร์ได้ ดูรายละเอียดเพิ่มเติมได้ที่ custom_call

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

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

ประเภทองค์ประกอบแสดงองค์ประกอบของประเภทเทนเซอร์ ประเภทเหล่านี้ไม่ใช่ประเภทระดับเฟิร์สคลาสใน StableHLO ซึ่งต่างจากภาษาโปรแกรมหลายภาษา ซึ่งหมายความว่าโปรแกรม StableHLO ไม่สามารถแสดงค่าของประเภทเหล่านี้ได้โดยตรง (ดังนั้นจึงเป็นเรื่องปกติที่จะแสดงค่าสเกลาร์ของประเภท 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 รวม
• ประเภททศนิยมอาจเป็นประเภทใดประเภทหนึ่งต่อไปนี้
  • f8E3M4, f8E4M3 และ f8E5M2 เป็นเลขทศนิยมแบบลอยตัว 8 บิตตาม ข้อกำหนด IEEE-754
  • f8E4M3FN และ f8E5M2 ที่สอดคล้องกับการเข้ารหัส E4M3 และ E5M2 ของรูปแบบ FP8 ตามลำดับที่อธิบายไว้ในรูปแบบ FP8 สำหรับ Deep Learning
  • f8E4M3FNUZ และ f8E5M2FNUZ ประเภทที่สอดคล้องกับการเข้ารหัส E4M3 และ E5M2 ของรูปแบบ FP8 ที่อธิบายไว้ใน รูปแบบตัวเลข 8 บิตสำหรับ Deep Neural Network
  • 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 (การปรับขนาดเล็ก) ที่อธิบายไว้ในข้อกำหนดรูปแบบการปรับขนาดเล็กของ OCP
• ประเภทเชิงซ้อนแสดงค่าเชิงซ้อนที่มีส่วนจริง และส่วนจินตภาพของประเภทองค์ประกอบเดียวกัน ประเภทที่ซับซ้อนที่รองรับคือ complex<f32> (ทั้ง 2 ส่วนเป็นประเภท f32) และ complex<f64> (ทั้ง 2 ส่วนเป็นประเภท f64)

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

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

StringType ::= 'string'

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

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

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

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

แอตทริบิวต์อินพุตมีชื่อและค่าซึ่งเป็นค่าคงที่ที่รองรับค่าใดค่าหนึ่ง ซึ่งเป็นวิธีหลักในการระบุข้อมูลเมตาแบบคงที่สำหรับองค์ประกอบของโปรแกรม ตัวอย่างเช่น concatenate op ใช้แอตทริบิวต์ 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}] ')'

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

ค่าคงที่จำนวนเต็มแสดงค่าจำนวนเต็มผ่านสตริงที่ใช้สัญกรณ์ทศนิยมหรือ เลขฐานสิบหก ระบบไม่รองรับฐานอื่นๆ เช่น ฐาน 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))
• (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

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

แบบสัมบูรณ์

ความหมาย

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

• สำหรับจำนวนเต็มที่มีเครื่องหมาย: โมดูลัสจำนวนเต็ม
• สำหรับ Float: 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]

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

เพิ่ม

ความหมาย

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

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

อินพุต

เอาต์พุต

ข้อจำกัด

• หากการดำเนินการใช้เทนเซอร์ที่ไม่ได้ทำ Quantize ให้ทำดังนี้
  • (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 จะต่อค่า ของเทนเซอร์ operands จากแต่ละกระบวนการตาม all_gather_dim และสร้างเทนเซอร์ results

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

• cross_replica(replica_groups) หาก channel_id <= 0 and use_global_device_ids = false
• cross_replica_and_partition(replica_groups) หาก channel_id > 0 and use_global_device_ids = false
• flattened_ids(replica_groups) หาก channel_id > 0 and use_global_device_ids = true

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

• operands...@receiver = [operand@sender for sender in process_group] สำหรับreceiverทั้งหมดใน process_group
• results...@process = concatenate(operands...@process, all_gather_dim) สำหรับprocessทั้งหมดใน process_group

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) 0 <= all_gather_dim < rank(operands...)
• (C2) is_unique(replica_groups)
• (C3) size(replica_groups) มีคำจำกัดความดังนี้
  • 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) หาก 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) มีคำจำกัดความดังนี้
  • 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...)
• (C7) element_type(results...) = E

ตัวอย่าง

// num_replicas: 2
// num_partitions: 1
// %operand0@(0, 0): [1, 2, 3, 4]
// %operand0@(1, 0): [5, 6, 7, 8]
// %operand1@(0, 0): [9, 10, 11, 12]
// %operand1@(1, 0): [13, 14, 15, 16]
%result:2 = "stablehlo.all_reduce"(%operand0, %operand0) ({
  ^bb0(%arg0: tensor<i64>, %arg1: tensor<i64>):
    %0 = "stablehlo.add"(%arg0, %arg1) : (tensor<i64>, tensor<i64>) -> tensor<i64>
    "stablehlo.return"(%0) : (tensor<i64>) -> ()
}) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  // channel_id = 0
  channel_handle = #stablehlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<4xi64>, tensor<4xi64>) -> (tensor<4xi64>, tensor<4xi64>)
// %result0@(0, 0): [6, 8, 10, 12]
// %result0@(1, 0): [6, 8, 10, 12]
// %result1@(0, 0): [22, 24, 26, 28]
// %result1@(1, 0): [22, 24, 26, 28]

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

all_to_all

ความหมาย

ในแต่ละกลุ่มกระบวนการในตารางกระบวนการ StableHLO จะแยกค่าของเทนเซอร์ operands ตาม split_dimension ออกเป็นส่วนๆ กระจายส่วนที่แยกแล้วระหว่างกระบวนการต่างๆ ต่อส่วนที่กระจายแล้วตาม concat_dimension และสร้างเทนเซอร์ results การดำเนินการจะแยกตารางกระบวนการ StableHLO ออกเป็น process_groups ซึ่งกำหนดไว้ดังนี้

• 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 รายการ lhs และ rhs แล้วสร้างเทนเซอร์ 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

ความหมาย

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

• สำหรับ Float: 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 โดยbatch_norm_trainingจะส่งต่อข้อผิดพลาดจาก grad_output และสร้างเทนเซอร์ grad_operand, grad_scale และ grad_offset ในรูปแบบที่เป็นทางการมากขึ้น การดำเนินการนี้สามารถแสดงเป็นการแยกย่อยเป็นการดำเนินการ StableHLO ที่มีอยู่โดยใช้ไวยากรณ์ Python ดังนี้

def compute_sum(operand, feature_index):
  (sum,) = reduce(
      inputs=[operand],
      init_values=[constant(0, element_type(operand))],
      dimensions=[i for i in range(rank(operand)) if i != feature_index],
      body=lambda x, y: add(x, y))
  return sum

def compute_mean(operand, feature_index):
  sum = compute_sum(operand, feature_index)
  divisor = constant(size(operand) / dim(operand, feature_index),
                     element_type(operand))
  divisor_bcast = broadcast_in_dim(divisor, [], type(sum))
  return divide(sum, divisor_bcast)

def batch_norm_grad(operand, scale, mean, variance, grad_output, epsilon, feature_index):
  # Broadcast inputs to type(operand)
  scale_bcast = broadcast_in_dim(scale, [feature_index], type(operand))
  mean_bcast = broadcast_in_dim(mean, [feature_index], type(operand))
  variance_bcast = broadcast_in_dim(variance, [feature_index], type(operand))
  epsilon_bcast = broadcast_in_dim(constant(epsilon, element_type(operand)), [],
                                   type(operand))

  # Perform normalization using the provided `mean` and `variance`
  # Intermediate values will be useful for computing gradients
  centered_operand = subtract(operand, mean_bcast)
  stddev = sqrt(add(variance_bcast, epsilon_bcast))
  normalized_operand = divide(centered_operand, stddev)

  # Use the implementation from batchnorm_expander.cc in XLA
  # Temporary variables have exactly the same names as in the C++ code
  elements_per_feature = broadcast_in_dim(
      constant(divide(size(operand), dim(operand, feature_index)),
               element_type(grad_output)),
      [], type(operand))
  i1 = multiply(grad_output, elements_per_feature)
  i2 = broadcast_in_dim(
      compute_sum(grad_output, feature_index), [feature_index], type(operand))
  i3 = broadcast_in_dim(
      compute_sum(multiply(grad_output, centered_operand), feature_index),
      [feature_index], type(operand))
  i4 = multiply(i3, centered_operand)
  i5 = divide(i4, add(variance_bcast, epsilon_bcast))
  i6 = subtract(subtract(i1, i2), i5)

  grad_operand =
      multiply(divide(divide(scale_bcast, stddev), elements_per_feature), i6)
  grad_scale =
      compute_sum(multiply(grad_output, normalized_operand), feature_index)
  grad_offset = compute_sum(grad_output, feature_index)

  return grad_operand, grad_scale, grad_offset

สำหรับประเภทที่กำหนดปริมาณ ให้ดำเนินการ dequantize_batch_norm_grad_or_training_quantize(lambda operand, scale, mean, variance, grad_output: batch_norm_grad(operand, scale, mean, variance, grad_output, epsilon, feature_index), operand, scale, mean, variance, grad_output, type(grad_operand), type(grad_scale), type(feature_index))

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) 0 <= feature_index < rank(operand)
• (C2) operand, scale, mean, variance, grad_output, grad_operand, grad_scale และ grad_offset มี baseline_element_type เดียวกัน
• (ค3) 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 ในรูปแบบที่เป็นทางการมากขึ้น การดำเนินการนี้สามารถแสดงเป็นการ แยกย่อยเป็นการดำเนินการ 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)
• (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
  • 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) ตามลำดับ
• (C2) 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]
//            ]
//          ]

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

เคส

ความหมาย

สร้างเอาต์พุตจากการเรียกใช้ฟังก์ชันเพียง 1 ฟังก์ชันจาก branches โดยขึ้นอยู่กับค่าของ 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]

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

cbrt

ความหมาย

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

• สำหรับ Float: 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 positive-definite ระบบจะไม่กำหนดลักษณะการทำงาน

สำหรับประเภทที่กำหนดปริมาณ ให้ดำเนินการ 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 ซึ่งกำหนดไว้ดังนี้

• 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 จะส่งค่าของเทนเซอร์ operand จากกระบวนการต้นทางไปยังกระบวนการเป้าหมาย และสร้างเทนเซอร์ 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
• (C2) 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]]

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

เปรียบเทียบ

ความหมาย

ทําการเปรียบเทียบแบบทีละองค์ประกอบของเทนเซอร์ lhs และ rhs ตาม comparison_direction และ compare_type และสร้างเทนเซอร์ result

ค่าของ comparison_direction และ compare_type มีความหมายดังนี้

สำหรับประเภทองค์ประกอบบูลีนและจำนวนเต็ม ให้ทำดังนี้

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

สำหรับประเภทองค์ประกอบทศนิยมที่มี compare_type = FLOAT ตัวดำเนินการจะใช้การดำเนินการ IEEE-754 ต่อไปนี้

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

สำหรับประเภทองค์ประกอบทศนิยมที่มี compare_type = TOTALORDER ตัวดำเนินการจะใช้การผสมผสานระหว่างการดำเนินการ totalOrder และ compareQuietEqual จาก IEEE-754

สำหรับประเภทองค์ประกอบที่ซับซ้อน ระบบจะทำการเปรียบเทียบตามพจนานุกรมของคู่ (real, imag) โดยใช้ comparison_direction และ compare_type ที่ระบุ การกำหนดลำดับให้กับจำนวนเชิงซ้อนเกี่ยวข้องกับความหมายที่น่าประหลาดใจ ดังนั้นในอนาคตเราจึงวางแผนที่จะยกเลิกการรองรับจำนวนเชิงซ้อน เมื่อ comparison_direction เป็น GE, GT, LE หรือ LT (#560)

สำหรับประเภทที่เล็กลง ให้ดำเนินการ dequantize_compare(lhs, rhs, comparison_direction)

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) baseline_element_type(lhs) = baseline_element_type(rhs)
• (C2) shape(lhs) = shape(rhs) = shape(result)
• (C3) compare_type มีคำจำกัดความดังนี้
  • 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)
• (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 op ด้วยการแยกย่อยโดยไม่เปลี่ยนความหมายของโปรแกรม ในกรณีที่การแทรกการแยกย่อยไม่ได้ให้ความหมายของ Op เดียวกัน ให้ใช้ 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>

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

เชื่อมโยง

ความหมาย

ต่อกัน 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

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

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

สำหรับ Conversion ที่เกี่ยวข้องกับ floating-point-to-integer ระบบจะตัดส่วนเศษส่วนออก หากค่าที่ถูกตัดทอนแสดงในประเภทปลายทางไม่ได้ ลักษณะการทำงานจะยังไม่กำหนด (#180)

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

สำหรับ Conversion complex-to-any-other-type และ any-other-type-to-complex ระบบจะไม่สนใจมูลค่าสมมติของแหล่งที่มาหรือตั้งค่ามูลค่าสมมติของปลายทางเป็น ศูนย์ตามลำดับ การแปลงส่วนจริงจะเป็นไปตาม การแปลงจุดลอย

ในทางทฤษฎี การดำเนินการนี้สามารถแสดงการยกเลิกการกำหนดปริมาณ (การแปลงจากเทนเซอร์ที่กำหนดปริมาณเป็นเทนเซอร์ปกติ) การกำหนดปริมาณ (การแปลงจากเทนเซอร์ปกติเป็นเทนเซอร์ที่กำหนดปริมาณ) และการกำหนดปริมาณใหม่ (การแปลงระหว่างเทนเซอร์ที่กำหนดปริมาณ) แต่ในขณะนี้เรามีการดำเนินการเฉพาะสำหรับการดำเนินการดังกล่าว - uniform_dequantize สำหรับกรณีการใช้งานแรก และ uniform_quantize สำหรับกรณีการใช้งานที่ 2 และ 3 ในอนาคต เราอาจรวม 2 การดำเนินการนี้ เข้ากับ convert (#1576)

อินพุต

เอาต์พุต

ข้อจำกัด

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

ตัวอย่าง

// %operand: [-1, 0, 1]
%result = "stablehlo.convert"(%operand) : (tensor<3xi64>) -> tensor<3xcomplex<f64>>
// %result: [(-1.0, 0.0), (0.0, 0.0), (1.0, 0.0)]

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

การสังวัตนาการ

ความหมาย

คำนวณผลคูณสเกลาร์ระหว่างหน้าต่างของ lhs กับชิ้นของ rhs และสร้าง result แผนภาพต่อไปนี้แสดงวิธีคำนวณองค์ประกอบใน result จาก lhs และ rhs โดยใช้ตัวอย่างที่เป็นรูปธรรม

กล่าวอย่างเป็นทางการมากขึ้น ให้พิจารณาการปรับเปลี่ยนอินพุตต่อไปนี้ในแง่ของ lhs เพื่อที่จะแสดงหน้าต่างของ lhs ได้

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

การปรับเฟรมนี้ใช้ฟังก์ชันตัวช่วยต่อไปนี้

• lhs_shape(n, hw, c) = permute([n] + hw + [c], [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension])
• result_shape(n1, hw, c1) = permute([n1] + hw + [c1], [output_batch_dimension] + output_spatial_dimensions + [output_feature_dimension])
• permute([j0, j1, ..., jR-1], permutation) = [i0, i1, ..., iR-1] เมื่อ j[d] = i[permutation[d]]

หาก feature_group_count = 1 และ batch_group_count = 1 ให้ทำดังนี้สำหรับทุก output_spatial_index ใน index_space(dim(result, output_spatial_dimensions...)) result[result_shape(:, output_spatial_index, :)] = dot_product โดยที่

• padding_value = constant(0, element_type(lhs))
• padded_lhs = pad(lhs, padding_value, lhs_padding[:, 0], lhs_padding[:, 1], lhs_base_dilations - 1)
• lhs_window_start = lhs_shape(0, output_spatial_index, 0) * lhs_window_strides
• lhs_window = slice(padded_lhs, lhs_window_start, lhs_window_start + lhs_window_dimensions, lhs_window_dilations)
• reversed_lhs_window = reverse(lhs_window, [input_spatial_dimensions[dim] for dim in range(size(window_reversal)) if window_reversal[dim] = true]) ดูเหมือนว่าไม่มีการใช้ฟีเจอร์นี้ ดังนั้นในอนาคตเราจึงวางแผนที่จะนำฟีเจอร์นี้ออก (#1181)
• dot_product = dot_general(reversed_lhs_window, rhs, lhs_batching_dimensions=[], lhs_contracting_dimensions=input_spatial_dimensions + [input_feature_dimension], rhs_batching_dimensions=[], rhs_contracting_dimensions=kernel_spatial_dimensions + [kernel_input_feature_dimension])

หาก feature_group_count > 1

• lhses = split(lhs, feature_group_count, input_feature_dimension)
• rhses = split(rhs, feature_group_count, kernel_output_feature_dimension)
• results... = convolution(lhses..., rhses..., ..., feature_group_count=1, ...)
• result = concatenate(results, output_feature_dimension)

หาก batch_group_count > 1

• lhses = split(lhs, batch_group_count, input_batch_dimension)
• rhses = split(rhs, batch_group_count, kernel_output_feature_dimension)
• results... = convolution(lhses..., rhses..., ..., batch_group_count=1, ...)
• result = concatenate(results, output_feature_dimension)

สำหรับประเภทที่กำหนดปริมาณ ให้ดำเนินการ dequantize_op_quantize( lambda lhs, rhs: convolution(lhs, rhs, window_strides, padding, lhs_dilation, rhs_dilation, window_reversal, input_batch_dimension, input_feature_dimension, input_spatial_dimensions, kernel_input_feature_dimension, kernel_output_feature_dimension, kernel_spatial_dimensions, output_batch_dimension, output_feature_dimension, output_spatial_dimensions, feature_group_count, batch_group_count, precision_config), lhs, rhs, type(result))

สำหรับประเภทที่เล็กลงแบบไฮบริด ให้ดำเนินการ hybrid_dequantize_then_op( lambda lhs, rhs: convolution(lhs, rhs, window_strides, padding, lhs_dilation, rhs_dilation, window_reversal, input_batch_dimension, input_feature_dimension, input_spatial_dimensions, kernel_input_feature_dimension, kernel_output_feature_dimension, kernel_spatial_dimensions, output_batch_dimension, output_feature_dimension, output_spatial_dimensions, feature_group_count, batch_group_count, precision_config), lhs, rhs)

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) N = rank(lhs) = rank(rhs)
• (C2) size(window_strides) = N - 2
• (C3) 0 < window_strides
• (C4) shape(padding) = [N - 2, 2]
• (C5) size(lhs_dilation) = N - 2
• (C6) 0 < lhs_dilation
• (C7) size(rhs_dilation) = N - 2
• (C8) 0 < rhs_dilation
• (C9) size(window_reversal) = N - 2
• (C10) dim(lhs, input_batch_dimension) % batch_group_count = 0
• (C11) dim(lhs, input_feature_dimension) % feature_group_count = 0
• (C12) size(input_spatial_dimensions) = N - 2
• (C13) เมื่อกำหนด input_dimensions = [input_batch_dimension] + input_spatial_dimensions + [input_feature_dimension]:
  • 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
• หากการดำเนินการใช้เทนเซอร์ที่ไม่ได้ทำ Quantize ให้ทำดังนี้
  • (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 โดยจะทำสิ่งต่อไปนี้ ทั้งนี้ขึ้นอยู่กับประเภทองค์ประกอบ

• สำหรับ Float: 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)

อินพุต

เอาต์พุต

(การรองรับ GPU ของ XLA) เป้าหมาย custom_call พิเศษ

มีฟังก์ชันพิเศษ 3 อย่างที่เกี่ยวข้องกับประเภท buffer ดังนี้ CreateBuffer สร้าง buffer ที่ยังไม่ได้เริ่มต้น Pin สร้าง buffer ที่เริ่มต้นแล้ว และ Unpin ยกเลิกการจัดสรร buffer และส่งคืนเนื้อหาของ buffercall_target_name

%uninitialized_buffer = "stablehlo.custom_call"() {
  call_target_name = "CreateBuffer",
  api_version = 4 : i32,
} : () -> memref<4xf64>

%initialized_buffer = "stablehlo.custom_call"(%init_value) {
  call_target_name = "Pin",
  api_version = 4 : i32,
} : (tensor<4xf64>) -> memref<4xf64>

%dealloc_buffer = "stablehlo.custom_call"(%initialized_buffer) {
  call_target_name = "Unpin",
  api_version = 4 : i32,
} : (memref<4xf64>) -> tensor<4xf64>

ชื่อแทน

การดำเนินการ custom_call บางอย่างอาจต้องใช้ส่วนหนึ่งในเอาต์พุตและอีกส่วนหนึ่งใน ตัวถูกดำเนินการเพื่อใช้หน่วยความจำเดียวกัน ซึ่งแสดงได้ผ่าน output_operand_aliases การแสดงคู่แทนประกอบด้วยรายการดัชนีเอาต์พุต ของ Tuple ที่แสดงส่วนเอาต์พุต และ operand_index พร้อมกับรายการดัชนี Tuple ของตัวถูกดำเนินการ ที่แสดงส่วนตัวถูกดำเนินการ รายการดัชนีเอาต์พุต หรือดัชนีทูเพิลตัวถูกดำเนินการจะว่างเปล่าหากประเภทที่เกี่ยวข้องไม่ใช่ประเภท tuple และอาจยาวเท่าใดก็ได้สำหรับประเภททูเพิลที่ซ้อนกันโดยพลการ ซึ่งคล้ายกับการแสดงแทนชื่อแทน XLA

ส่วนเอาต์พุตและส่วนอินพุตในคู่แทนต้องมีประเภทเดียวกัน สำหรับ การดำเนินการ custom_call ที่ไม่ใช่การเรียกไปยัง CreateBuffer, Pin และ Unpin ตัว bufferตัวถูกดำเนินการจะปรากฏในคู่ชื่อแทนได้สูงสุด 1 คู่ และเอาต์พุต buffer ต้องปรากฏในคู่ชื่อแทน 1 คู่

ตัวอย่าง

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

%updated_buffer = "stablehlo.custom_call"(%buffer) {
  call_target_name = "Update",
  api_version = 4 : i32,
  output_operand_aliases = [
    #stablehlo.output_operand_alias<output_tuple_indices = [],
      operand_index = 0,
      operand_tuple_indices = []>]
} : (memref<4xf64>) -> memref<4xf64>

หาร

ความหมาย

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

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

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

อินพุต

เอาต์พุต

ข้อจำกัด

• (C1) element_type(result) มีค่าดังนี้
  • element_type(operand) หาก !is_per_axis_quantized(operand)
  • element_type(operand) ยกเว้น quantization_dimension(operand), scales(operand) และ zero_points(operand) อาจแตกต่างจาก quantization_dimension(result), scales(result) และ zero_points(result) ตามลำดับ
• (C2) size(broadcast_dimensions) = rank(operand)
• (C3) 0 <= broadcast_dimensions < rank(result)
• (C4) is_unique(broadcast_dimensions)
• (C5) สำหรับ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

ความหมาย

การดำเนินการนี้จะเหมือนกับ การแปลงคอนโวลูชัน ในเชิงฟังก์ชัน แต่จะมีการระบุการเพิ่มพื้นที่แบบไดนามิกผ่าน padding

อินพุต