'мхло' Диалект

Операции

mhlo.abs (mhlo::AbsOp)

Операция на прессе

Синтаксис:

operation ::= `mhlo.abs` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Выполняет поэлементную операцию абс над тензором operand и создает тензор result .

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#abs .

Пример:

%result = mhlo.abs %operand : tensor<3xi32>

Признаки: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape .

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
operand ранговый тензор 2/4/8/16/32/64-битного целого числа без знака или типа f4E2M1FN, или типа f6E2M3FN, или типа f6E3M2FN, или типа f8E3M4, или типа f8E4M3, или типа f8E4M3FN, или типа f8E4M3FNUZ, или типа f8E4M3B11FNUZ, или типа f8E5M2, или f8E5M2FNU Тип Z или f8E8M0FNU тип или 16-битное число с плавающей запятой, или 32-битное число с плавающей запятой, или 64-битное число с плавающей запятой, или тип bfloat16 или комплексный тип с 32-битным числом с плавающей запятой или 64-битное число с плавающей запятой, или 2/4/8/16/32-битное равномерное квантованное целое число со знаком или 2/4/8/16/32-битное целое число с равномерным квантованием по оси со знаком или 2/4/8/16/32-битное целое число с равномерным квантованием без знака или 2/4/8/16/32-битное целое число с равномерным квантованием по оси без знака целочисленные значения

Результаты:

Результат Описание
result ранговый тензор 2/4/8/16/32/64-битного целого числа без знака или типа f4E2M1FN, или типа f6E2M3FN, или типа f6E3M2FN, или типа f8E3M4, или типа f8E4M3, или типа f8E4M3FN, или типа f8E4M3FNUZ, или типа f8E4M3B11FNUZ, или типа f8E5M2, или f8E5M2FNU Тип Z или f8E8M0FNU тип или 16-битное число с плавающей запятой, или 32-битное число с плавающей запятой, или 64-битное число с плавающей запятой, или тип bfloat16, или 2/4/8/16/32-битное равномерно квантованное целое число со знаком или 2/4/8/16/32-битное равномерно квантованное число на каждый Целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака или 2/4/8/16/32-битное равномерно квантованное целое число без знака для каждой оси

mhlo.add (mhlo::AddOp)

Добавить операцию

Синтаксис:

operation ::= `mhlo.add` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Выполняет поэлементное сложение двух тензоров lhs и rhs и создает result тензор.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#add .

Пример:

%result = mhlo.add %lhs, %rhs : tensor<2x2xi32>

Признаки: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
lhs ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битные элементы с плавающей точкой или 64-битные элементы с плавающей запятой или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или 2/4/8/16 /32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по каждой оси целое число без знака
rhs ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битные элементы с плавающей точкой или 64-битные элементы с плавающей запятой или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или 2/4/8/16 /32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по оси целое число без знака

Результаты:

Результат Описание
result ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битные элементы с плавающей точкой или 64-битные элементы с плавающей запятой или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или 2/4/8/16 /32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по каждой оси целое число без знака

mhlo.add_dependency (mhlo::AddDependencyOp)

Операция AddDependency

Синтаксис:

operation ::= `mhlo.add_dependency` operands attr-dict `:` functional-type(operands, results)

Эта операция является частной для компилятора XLA, поэтому для нее еще нет спецификации.

Неформально у этой операции два операнда: операнд данных и токен. Результатом операции является операнд данных. При использовании с AfterAll эта операция позволяет упорядочивать операции без побочных эффектов (те, которые не создают значения токенов).

Пример:

%1 = mhlo.add_dependency %arg0, %0 : (tensor<3x4xf32>, !mhlo.token) -> tensor<3x4xf32>

Черты: AlwaysSpeculatableImplTrait

Интерфейсы: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
operand ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битные элементы с плавающей запятой или 64-битные элементы с плавающей запятой или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или ранговый тензор 2/4 /8/16/32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по каждой оси целое число без знака или токен
token жетон

Результаты:

Результат Описание
output ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битные элементы с плавающей запятой или 64-битные элементы с плавающей запятой или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или ранговый тензор 2/4 /8/16/32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по каждой оси целое число без знака или токен

mhlo.after_all (mhlo::AfterAllOp)

После операции

Синтаксис:

operation ::= `mhlo.after_all` $inputs attr-dict
              `:` custom<VariadicSameOperandsAndResultType>(ref($inputs), type($inputs), type($result))

Гарантирует, что операции, производящие inputs , выполняются до выполнения любых операций, зависящих от result .

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all.

Пример:

%result = mhlo.after_all %input0, %input1 : !mhlo.token

Черты: AlwaysSpeculatableImplTrait

Интерфейсы: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
inputs вариант токена

Результаты:

Результат Описание
result жетон

mhlo.all_gather (mhlo::AllGatherOp)

Операция AllGather

Внутри каждой группы процессов в сетке процессов объединяет значения тензора операндов каждого процесса по all_gather_dim и создает результирующий тензор. computation применяются отдельно для каждого операнда в operands , давая один результат для каждого операнда.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_gather .

Пример:

%result = "mhlo.all_gather"(%operand) {
  all_gather_dim = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>,
  // use_global_device_ids = false
} : (tensor<2x2xf32>) -> tensor<2x4xf32>

Признаки: SameOperandsAndResultElementType

Атрибуты:

Атрибут Тип МЛИР Описание
all_gather_dim ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого неотрицательно.
replica_groups ::mlir::DenseIntElementsAttr Атрибут 64-битного целочисленного элемента без знака
channel_handle ::mlir::mhlo::ChannelHandleAttr два 64-битных целых числа «дескриптор» и «тип»
use_global_device_ids ::mlir::UnitAttr атрибут единицы

Операнды:

Операнд Описание
operands вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси

Результаты:

Результат Описание
«безымянный» вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей точкой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси

mhlo.all_reduce (mhlo::AllReduceOp)

Операция AllReduce

В каждой группе процессов в сетке процессов применяет computation функции сокращения к значениям тензора операнда каждого процесса и создает результирующий тензор. computation применяются отдельно для каждого операнда в operands , давая один результат для каждого операнда.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce.

Пример:

%result = "mhlo.all_reduce"(%operand) ({
  ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>):
    %0 = mhlo.add %arg1, %arg2 : tensor<f32>
    mhlo.return %0 : tensor<f32>
}) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<4xf32>) -> tensor<4xf32>

Признаки: InferTensorType , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Интерфейсы: InferShapedTypeOpInterface , InferTypeOpInterface

Атрибуты:

Атрибут Тип МЛИР Описание
replica_groups ::mlir::DenseIntElementsAttr Атрибут 64-битного целочисленного элемента без знака
channel_handle ::mlir::mhlo::ChannelHandleAttr два 64-битных целых числа «дескриптор» и «тип»
use_global_device_ids ::mlir::UnitAttr атрибут единицы

Операнды:

Операнд Описание
operands вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей точкой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси

Результаты:

Результат Описание
«безымянный» вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей точкой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси

mhlo.all_to_all (mhlo::AllToAllOp)

Операция AllToAll

Внутри каждой группы процессов в сетке процессов разбивает значения тензора operand по split_dimension на части, распределяет разделенные части между процессами, объединяет разбросанные части по concat_dimension и создает result тензор.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_to_all .

Пример:

%result = "mhlo.all_to_all"(%operand) {
  split_dimension = 1 : i64,
  concat_dimension = 0 : i64,
  split_count = 2 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
} : (tensor<2x4xf32>) -> tensor<4x2xf32>

Признаки: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsElementType , SameOperandsShape , SameVariadicOperandSize

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип МЛИР Описание
split_dimension ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого неотрицательно.
concat_dimension ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого неотрицательно.
split_count ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого положительное.
replica_groups ::mlir::DenseIntElementsAttr Атрибут 64-битного целочисленного элемента без знака
channel_handle ::mlir::mhlo::ChannelHandleAttr два 64-битных целых числа «дескриптор» и «тип»

Операнды:

Операнд Описание
operand вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси

Результаты:

Результат Описание
«безымянный» вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей точкой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси

mhlo.and (mhlo::AndOp)

И операция

Синтаксис:

operation ::= `mhlo.and` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Выполняет поэлементное И двух тензоров lhs и rhs и создает result тензор.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#and

Пример:

%result = mhlo.and %lhs, %rhs : tensor<2x2xi32>

Признаки: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
lhs ранжированный тензор pred (также известный как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака.
rhs ранжированный тензор pred (также известный как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака.

Результаты:

Результат Описание
result ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битные элементы с плавающей запятой или 64-битные элементы с плавающей точкой или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или 2/4/8/16 /32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по оси целое число без знака

mhlo.async_done (mhlo::AsyncDoneOp)

Операция AsyncDone

Эта операция является частной для компилятора XLA, поэтому для нее еще нет спецификации.

Неформально эта операция блокируется до конца асинхронных вычислений. Он возвращает конечный результат асинхронных вычислений.

Дополнительную информацию см. в документации по AsyncStart.

Интерфейсы: InferTypeOpInterface

Операнды:

Операнд Описание
bundle async_bundle с любой комбинацией ранжированного тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или f8E8M0FNU тип или 16-битное число с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой, или тип bfloat16, или pred (также известное как логическое или 1-битное целое число), или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака. Целочисленный или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей точкой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/ 4/8/16/32-битные, равномерно квантованные по каждой оси целые числа со знаком или 2/4/8/16/32-битные, равномерно квантованные по каждой оси целые числа без знака или значения маркеров

Результаты:

Результат Описание
«безымянный» вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное целое число со знаком по оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака по каждой оси или токен или вложенный кортеж с любой комбинацией ранжированного тензора типа f4E2M1FN или типа f6E2M3FN или f6E3M2FN тип или тип f8E3M4 или тип f8E4M3 или тип f8E4M3FN или тип f8E4M3FNUZ или тип f8E4M3B11FNUZ или тип f8E5M2 или тип f8E5M2FNUZ или тип f8E8M0FNU или 16-битное число с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (AKA boolean или 1) -битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битными или 64-битными элементами с плавающей запятой. или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или ранжированный тензор 2/4/8/16/32-битных равномерно квантованных чисел целое число со знаком на каждой оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака на каждой оси или значения токена

mhlo.async_start (mhlo::AsyncStartOp)

Операция асинхронного запуска

Эта операция является частной для компилятора XLA, поэтому для нее еще нет спецификации.

Неформально эта операция запускает асинхронные вычисления.

Это используется, когда существуют функции, которые содержат как асинхронное ожидание (например, DMA), так и вычисления в потоке. Например, функция может состоять из вычисления, прямого доступа к памяти, другого вычисления, второго прямого доступа к памяти и окончательного вычисления. Это будет представлено как async_start, за которым следуют async_update и async_done. async_start выполнит первые вычисления в потоке, а затем запустит DMA. async_update будет ждать завершения DMA, если оно еще не было выполнено, затем выполнит второе вычисление в функции и запустит второй DMA. Наконец, async_done будет ждать этого последнего DMA, а затем запускает последнее вычисление, которое необходимо запустить в потоке, и возвращает результат этого окончательного вычисления.

operands передаются в вычисление напрямую. called_computation — это функция, которая будет выполняться асинхронно. execution_thread — это имя потока, в котором она будет выполняться. Основной поток называется «главным». Все темы имеют названия.

Это возвращает все состояние, необходимое между асинхронными операциями. После назначения буфера возвращаемые значения представляют собой пространство, необходимое для хранения входных данных, результатов и любых блокнотов, необходимых или редактируемых асинхронной операцией.

Атрибуты:

Атрибут Тип МЛИР Описание
called_computation ::mlir::FlatSymbolRefAttr Атрибут ссылки на плоский символ
execution_thread ::mlir::StringAttr строковый атрибут

Операнды:

Операнд Описание
inputs вариация рангового тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16- бит с плавающей запятой или 32-битный с плавающей запятой или 64-битный бит с плавающей запятой или тип bfloat16 или pred (также известное как логическое или 1-битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексное число тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/4/8 /16/32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по оси целое число без знака или токен или вложенный кортеж с любой комбинацией ранжированного тензора типа f4E2M1FN или типа f6E2M3FN или f6E3M2FN тип или тип f8E3M4 или тип f8E4M3 или тип f8E4M3FN или тип f8E4M3FNUZ или тип f8E4M3B11FNUZ или тип f8E5M2 или тип f8E5M2FNUZ или тип f8E8M0FNU или 16-битное число с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или pred (AKA boolean или 1) -битное целое число) или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака или комплексный тип с 32-битными или 64-битными элементами с плавающей запятой. или 2/4/8/16/32-битные равномерно квантованные целые числа со знаком или 2/4/8/16/32-битные равномерно квантованные целые числа без знака или ранжированный тензор 2/4/8/16/32-битных равномерно квантованных чисел целое число со знаком на каждой оси или 2/4/8/16/32-битное равномерно квантованное целое число без знака на каждой оси или значения токена

Результаты:

Результат Описание
«безымянный» async_bundle с любой комбинацией ранжированного тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или f8E8M0FNU тип или 16-битное число с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой, или тип bfloat16, или pred (также известное как логическое или 1-битное целое число), или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака. Целочисленный или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей точкой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/ 4/8/16/32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по каждой оси целое число без знака или значения токена

mhlo.async_update (mhlo::AsyncUpdateOp)

Операция AsyncUpdate

Эта операция является частной для компилятора XLA, поэтому для нее еще нет спецификации.

Неформально эта операция блокирует асинхронные вычисления до тех пор, пока не возникнет барьер синхронизации. Это возвращает bundle после работы с ним.

Дополнительную информацию см. в документации по AsyncStart.

Interfaces: InferTypeOpInterface

Операнды:

Операнд Описание
bundle async_bundle с любой комбинацией ранжированного тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или f8E8M0FNU тип или 16-битное число с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой, или тип bfloat16, или pred (также известное как логическое или 1-битное целое число), или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака. Целочисленный или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей точкой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/ 4/8/16/32-битные, равномерно квантованные по каждой оси целые числа со знаком или 2/4/8/16/32-битные, равномерно квантованные по каждой оси целые числа без знака или значения маркеров

Результаты:

Результат Описание
«безымянный» async_bundle с любой комбинацией ранжированного тензора типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или f8E8M0FNU тип или 16-битное число с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой, или тип bfloat16, или pred (также известное как логическое или 1-битное целое число), или 2/4/8/16/32/64-битное целое число без знака или 2/4/8/16/32/64-битное целое число без знака. Целочисленный или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей точкой или 2/4/8/16/32-битным равномерным квантованным целым числом со знаком или 2/4/8/16/32-битным равномерным квантованным целым числом без знака или 2/ 4/8/16/32-битное равномерно квантованное по оси целое число со знаком или 2/4/8/16/32-битное равномерно квантованное по каждой оси целое число без знака или значения токена

mhlo.atan2 (mhlo::Atan2Op)

Операция Атан2

Синтаксис:

operation ::= `mhlo.atan2` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Выполняет поэлементную операцию atan2 над тензорами lhs и rhs и создает result тензор.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#atan2 .

Пример:

%result = mhlo.atan2 %lhs, %rhs : tensor<3xf32>

Признаки: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
lhs ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16 или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой, или 2/4/8/16/32-битными равномерными квантованными целыми числами со знаком или 2/4/8/16/32-битными равномерными квантованными целыми числами без знака.
rhs ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16, или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой, или 2/4/8/16/32-битное равномерно квантованное целое число со знаком или 2/4/8/16/32-битное равномерно квантованное целое число без знака.

Результаты:

Результат Описание
result ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или тип bfloat16, или комплексный тип с 32-битными элементами с плавающей запятой или 64-битными элементами с плавающей запятой, или 2/4/8/16/32-битное равномерно квантованное целое число со знаком или 2/4/8/16/32-битное равномерно квантованное целое число без знака.

mhlo.batch_norm_grad (mhlo::BatchNormGradOp)

Операция «БатчНормГрад»

Вычисляет градиенты нескольких входных данных BatchNormTrainingOp с обратным распространением ошибки от grad_output и создает тензоры grad_operand , grad_scale и grad_offset .

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_grad .

Пример:

%grad_operand, %grad_scale, %grad_offset =
"mhlo.batch_norm_grad"(%operand, %scale, %mean, %variance, %grad_output) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>,
    tensor<2x2x2xf32>) -> (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>)

Черты: AlwaysSpeculatableImplTrait , InferTensorType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип МЛИР Описание
epsilon ::mlir::FloatAttr 32-битный атрибут с плавающей запятой
feature_index ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого неотрицательно.

Операнды:

Операнд Описание
operand ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
scale 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
mean 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
variance 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
grad_output ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16

Результаты:

Результат Описание
grad_operand ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
grad_scale 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
grad_offset 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16

mhlo.batch_norm_inference (mhlo::BatchNormInferenceOp)

Операция BatchNormInference

Нормализует тензор operand по всем измерениям, кроме измерения feature_index , и создает result тензор.

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_inference .

Пример:

%result = "mhlo.batch_norm_inference"(%operand, %scale, %offset, %mean, %variance) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> tensor<2x2x2xf32>

Черты: AlwaysSpeculatableImplTrait , InferTensorType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип МЛИР Описание
epsilon ::mlir::FloatAttr 32-битный атрибут с плавающей запятой
feature_index ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого неотрицательно.

Операнды:

Операнд Описание
operand ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
scale 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
offset 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
mean 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16
variance 1D тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16

Результаты:

Результат Описание
result ранговый тензор типа f4E2M1FN или типа f6E2M3FN или типа f6E3M2FN или типа f8E3M4 или типа f8E4M3 или типа f8E4M3FN или типа f8E4M3FNUZ или типа f8E4M3B11FNUZ или типа f8E5M2 или типа f8E5M2FNUZ или типа f8E8M0FNU или 16-битного числа с плавающей запятой или 32-битное число с плавающей запятой или 64-битное число с плавающей запятой или значения типа bfloat16

mhlo.batch_norm_training (mhlo::BatchNormTrainingOp)

Пакетная нормаОбучающая операция

Вычисляет среднее значение и дисперсию по пакетам и пространственным измерениям, нормализует тензор operand для каждого объекта в измерении feature_index и создает output , тензоры batch_mean и batch_var .

См.: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_training .

Пример:

%output, %batch_mean, %batch_var = "mhlo.batch_norm_training"(%operand, %scale, %offset) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>) -> (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>)

Черты: AlwaysSpeculatableImplTrait , InferTensorType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип МЛИР Описание
epsilon ::mlir::FloatAttr 32-битный атрибут с плавающей запятой
feature_index ::mlir::IntegerAttr 64-битный целочисленный атрибут без знака, значение которого неотрицательно.

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или значения типа BFLOAT16
scale 1D Тензор типа F4E2M1FN или типа F6E2M3FN или тип f6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или значения типа BFLOAT16
offset 1D Тензор типа F4E2M1FN или типа F6E2M3FN или тип f6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или значения типа BFLOAT16

Результаты:

Результат Описание
output Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или значения типа BFLOAT16
batch_mean 1D Тензор типа F4E2M1FN или типа F6E2M3FN или тип f6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или значения типа BFLOAT16
batch_var 1D Тензор типа F4E2M1FN или типа F6E2M3FN или тип f6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или значения типа BFLOAT16

mhlo.bitcast (mhlo :: bitcastop)

Операция биткала

Синтаксис:

operation ::= `mhlo.bitcast` operands attr-dict `:` functional-type(operands, results)

Эта операция является частной для компилятора XLA, поэтому она еще не имеет спецификации.

Неофициально, эта операция изменяет форму ввода в том, как физическое расположение элементов не изменилось.

Эта операция нуждается в информации макета, чтобы понять «физическое расположение элементов», а поддержка макета в MHLO в настоящее время находится в стадии разработки.

Пример:

%0 = mhlo.bitcast %arg0 : (tensor<3x4xf32>) -> tensor<3x4x1xf32>

Черты: AlwaysSpeculatableImplTrait

Интерфейсы: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.bitcast_convert (mhlo :: bitcastconvertop)

Операция BitCastConvert

Синтаксис:

operation ::= `mhlo.bitcast_convert` operands attr-dict `:` functional-type(operands, results)

Выполняет операцию Bitcast на тензоре operand и дает тензор result , где биты всего тензора operand переосмысливаются с использованием типа тензора result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#bitcast_convert

Пример:

%result = mhlo.bitcast_convert %operand : (tensor<2xf32>) -> tensor<2x4xi8>

Черты: AlwaysSpeculatableImplTrait

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.broadcast (mhlo :: froadcastop)

Операция вещания

Эта операция выходит из stablehlo, поэтому она не включена в спецификацию: https://github.com/openxla/stablehlo/issues/3

Неофициально, эта операция делает то же самое, что и трансляция XLA: https://www.tensorflow.org/xla/operation_semantics#broadcast

Пример:

%result = mhlo.broadcast %operand, sizes = [1, 2] : (tensor<3xi32>) -> tensor<1x2x3xi32>

Черты: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
broadcast_sizes :: mlir :: denseintelementsattr 64-битный атрибут без знаковых целостных элементов

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.broadcast_in_dim (mhlo :: broadcastNidimop)

РАЗВИТИЯ ВВЕДЕНИЯ

Расширяет размеры и/или ранг входного тензора, дублируя данные в тензоре operand и дает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#broadcast_in_dim

Пример:

%result = mhlo.broadcast_in_dim %operand, dims = [2, 1] : (tensor<1x3xi32>) -> tensor<2x3x2xi32>

Черты: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Интерфейсы: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
broadcast_dimensions :: mlir :: denseintelementsattr 64-битный атрибут без знаковых целостных элементов

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Статическая форма тензора типа F4E2M1FN или типа F6E2M3FN или тип F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E5M3FNUZ или тип F8E-2 он плавает или 32-битный поплавок или 64-битный Float или Bfloat16 Тип или Pred (AKA Boolean или 1-битное целое число) или 2/4/8/16/32/64-битного. С 32-разрядными плавуческими или 64-разрядными плавучими элементами или 2/4/8/16/32-битным равномерным квантованным целым числом или 2/4/8/16/32-битом равномерного квантового квантового целого числа или 2/4/8// 16/32-битный равномерный квантовый на оси.

mhlo.case (mhlo :: caseop)

Дело

Производит выход из выполнения ровно одной function из branches в зависимости от значения index .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#case

Пример:

%result0, %result1 = "mhlo.case"(%index) ({
  mhlo.return %result_branch0, %result_branch0 : tensor<2xi64>, tensor<2xi64>
}, {
  mhlo.return %result_branch1, %result_branch1 : tensor<2xi64>, tensor<2xi64>
}) : (tensor<i32>) -> (tensor<2xi64>, tensor<2xi64>)

Черты: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Интерфейсы: InferTypeOpInterface

Операнды:

Операнд Описание
index Тензор 32-разрядных значений без знаков.

Результаты:

Результат Описание
«Безымянный» Вариада по ранжированию типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или типа F8E3M4 или типа F8E4M3B11FNE8 16-битный поплавок или 32-битный поплавок или 64- битовая плавание или тип Bfloat16 или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-разрядочное неверное целое число или 2/4/8/16/32/64-разрядное целое число или комплекс Введите с 32-разрядными платными или 64-битными поплавковыми элементами или 2/4/8/16/32-разрядным равномерным квантованным подписанным целым числом или 2/4/8/16/32-разрядом, равномерными квантованными квантованными целыми значениями без знака или по ранжированному тензору 2 /4/8/16/32-битный равномерный квантовый квантовый на оси целое число или 2/4/8/16/32-бита.

mhlo.cbrt (mhlo :: cbrtop)

Операция CBRT

Синтаксис:

operation ::= `mhlo.cbrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Выполняет элементную работу кубического корня на тензоре operand и получает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cbrt

Пример:

%result = mhlo.cbrt %operand : tensor<4xf32>

Черты: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или Bfloat16 или комплексный тип с 32-разрядными платными или 64-битными поплавковыми элементами или 2/4/8/16/32-битом равномерного квантового знакомого целого числа или 2/4/8/16/32-битных равномерных квантовых квантовых целочисленных значений без знака

Результаты:

Результат Описание
result Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или Bfloat16 или комплексный тип с 32-разрядными платными или 64-битными поплавковыми элементами или 2/4/8/16/32-битом равномерного квантового знакомого целого числа или 2/4/8/16/32-битных равномерных квантовых квантовых целочисленных значений без знака

mhlo.ceil (mhlo :: ceilop)

CEIL Operation

Синтаксис:

operation ::= `mhlo.ceil` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Выполняет элементный Ceil из operand Tensor и дает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#ceil

Пример:

%result = mhlo.ceil %operand : tensor<5xf32>

Черты: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или BFLOAT16 Тип или 2/4/8/16/32-битного равномерного квантового знакового целого числа или 2/4/8/16/32-битных равномерных квантовых целочисленных значений без знака

Результаты:

Результат Описание
result Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или BFLOAT16 Тип или 2/4/8/16/32-битного равномерного квантового знакового целого числа или 2/4/8/16/32-битных равномерных квантовых целочисленных значений без знака

mhlo.cholesky (mhlo :: chouleskyop)

Чолская операция

Вычисляет размесщение хоузского партии матриц.

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cholesky

Пример:

%result = mhlo.cholesky %a, lower = true : tensor<3x3xf32>

Черты: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
lower :: mlir :: boolattr атрибут Bool

Операнды:

Операнд Описание
a Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или тип Bfloat16 или комплексный тип с 32-разрядными значениями плавания или 64-битных поплавковых элементов

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или тип Bfloat16 или комплексный тип с 32-разрядными значениями плавания или 64-битных поплавковых элементов

mhlo.clamp (mhlo::ClampOp)

Операция зажима

Синтаксис:

operation ::= `mhlo.clamp` $min `,` $operand `,` $max attr-dict
              `:` custom<SameOperandsAndResultType>(type($min), type($operand), type($max), type($result))

Зажимая каждый элемент тензора operand между минимальным и максимальным значением и дает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#clamp

Пример:

%result = mhlo.clamp %min, %operand, %max : tensor<3xi32>

Черты: AlwaysSpeculatableImplTrait , HLO_BroadcastingElementwise , InferTensorType , SameOperandsAndResultElementType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
min Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения
max Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
result Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.collective_broadcast (mhlo :: collectivebroadcastop)

CollectiveBroadcast Operation

В каждой группе процессов в сетке процесса отправьте значение тензора operand из исходного процесса в целевые процессы и дают тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_broadcast

Пример:

%result = "mhlo.collective_broadcast"(%operand) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
} : (tensor<1x2xi64>) -> tensor<1x2xi64>

Черты: CompatibleOperandsAndResultType

Интерфейсы: InferShapedTypeOpInterface , InferTypeOpInterface

Атрибуты:

Атрибут Тип Mlir Описание
replica_groups :: mlir :: denseintelementsattr 64-битный атрибут без знаковых целостных элементов
channel_handle :: mlir :: mhlo :: cannelhandleattr Два 64-битных целых числа «ручка» и «тип»

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.collective_permute (mhlo :: collectivepermuteop)

CollectivePermute Operate

В каждой группе процессов в сетке процесса отправляет значение тензора operand из исходного процесса в целевой процесс и дает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_permute

Пример:

%result = "mhlo.collective_permute"(%operand) {
  source_target_pairs = dense<[[0, 1], [1, 2]]> : tensor<2x2xi64>,
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
} : (tensor<4x2xf32>) -> tensor<4x2xf32>

Черты: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
source_target_pairs :: mlir :: denseintelementsattr 64-битный атрибут без знаковых целостных элементов
channel_handle :: mlir :: mhlo :: cannelhandleattr Два 64-битных целых числа «ручка» и «тип»

Операнды:

Операнд Описание
operand Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.compare (mhlo :: compareop)

Сравните операцию

Синтаксис:

operation ::= `mhlo.compare` $comparison_direction `,` $lhs `,` $rhs (`,` $compare_type^)?
              attr-dict `:` functional-type(operands, results)

Проводит элементное сравнение тензоров lhs и rhs в соответствии с comparison_direction и compare_type и дает result тензора.

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#compare

Пример:

%result = mhlo.compare LT, %lhs, %rhs, FLOAT : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>

Черты: AlwaysSpeculatableImplTrait , Elementwise , InferTensorType , SameOperandsAndResultShape , SameOperandsElementType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
comparison_direction :: mlir :: mhlo :: compassondirectionattr Какую операцию сравнения выполнить.
compare_type :: mlir :: mhlo :: comparrishoniseTypeattr Какой тип сравнения использовать.

Операнды:

Операнд Описание
lhs Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения
rhs Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

Результаты:

Результат Описание
«Безымянный» Ранговый тензор значений Pred (AKA Boolean или 1-битного целочисленного)

mhlo.complex (mhlo :: complexop)

Сложная операция

Синтаксис:

operation ::= `mhlo.complex` operands attr-dict
              `:` custom<ComplexOpType>(type($lhs), type($rhs), type($result))

Выполняет элементное преобразование в сложное значение из пары реальных и воображаемых значений, lhs и rhs , и дает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#complex

Пример:

%result = mhlo.complex %lhs, %rhs : tensor<2xcomplex<f32>>

Черты: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape SameOperandsElementType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Операнды:

Операнд Описание
lhs Ранг Тензор 32-битного поплавкового или 64-битного поплавок значений
rhs Ранг Тензор 32-битного поплавкового или 64-битного поплавок значений

Результаты:

Результат Описание
result Ранг Тензор сложного типа с 32-разрядными значениями плавания или 64-битных поплавковых элементов

mhlo.composite (mhlo :: compositeop)

Композитная операция

Синтаксис:

operation ::= `mhlo.composite` $name $inputs attr-dict `:` functional-type(operands, results)

Инкапсулирует операцию, выполненную (составленную) других операций StableHlo, принимая inputs и composite_attributes и дает results . Семантика OP реализована атрибутом decomposition . composite OP может быть заменен его разложением без изменения семантики программы. В тех случаях, когда внедрение разложения не обеспечивает ту же операционную семантику, предпочитаю использовать custom_call .

Поле version (по умолчанию к 0 ) используется для обозначения при изменении семантики композита.

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#composite

Пример:

%results = mhlo.composite "my.op" %arg0, %arg1 {
  decomposition = @my_op,
  composite_attributes = { my_attribute = "my_value" },
  version = 1 : i32
} : (tensor<f32>, tensor<f32>) -> tensor<f32>

Интерфейсы: SymbolUserOpInterface

Атрибуты:

Атрибут Тип Mlir Описание
name :: mlir :: stringattr Строка атрибут
composite_attributes :: mlir :: dictionaryattr Словарь названных значений атрибутов
decomposition :: mlir :: flatsymbolrefattr Справочный атрибут плоского символа
version :: mlir :: integerattr 32-разрядный целочисленный атрибут.

Операнды:

Операнд Описание
inputs Вариада по ранжированию типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или типа F8E3M4 или типа F8E4M3B11FNE8 16-битный поплавок или 32-битный поплавок или 64- битовая плавание или тип Bfloat16 или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-разрядочное неверное целое число или 2/4/8/16/32/64-разрядное целое число или комплекс Введите с 32-разрядными платными или 64-битными поплавковыми элементами или 2/4/8/16/32-битным равномерным квантовым подписанным целым числом или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8 /16/32-битный равномерный квантованный на оси, подписанный целым числом или 2/4/8/16/32-битный, равномерный квантовый на оси не знагленные целочисленные значения или токен или вложенный трюк с любым комбинацией рапольного тензора F4E2M1FN или типа F6E2M3FN или F6E3M2FN Тип или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FN или тип F8E4M3FNUZ или тип F8E4M3B11FNUZ или тип F8E5M2 или F8E5M2FNUZ или тип F8E8M0FNU или 16-битный плавучий или 32-битный плавучий или 64-битный плавучий или BFLOAT16-тип. -bit integer) или 2/4/8/16/32/64-битный неверный целый integer или 2/4/8/16/32/64-разрядный не знаковый целое число или сложный тип с 32-битным поплавковым или 64-битным поплавковым элементом или 2/4/8/16/32-битный равномерный квантовый квантовый целое число или 2/4/8/16/32-бита. на оси, подписанное целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаменитые значения целочисленных или значений токена

Результаты:

Результат Описание
«Безымянный» Вариада по ранжированию типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или типа F8E3M4 или типа F8E4M3B11FNE8 16-битный поплавок или 32-битный поплавок или 64- битовая плавание или тип Bfloat16 или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-разрядочное неверное целое число или 2/4/8/16/32/64-разрядное целое число или комплекс Введите с 32-разрядными платными или 64-битными поплавковыми элементами или 2/4/8/16/32-битным равномерным квантовым подписанным целым числом или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8 /16/32-битный равномерный квантованный на оси, подписанный целым числом или 2/4/8/16/32-битный, равномерный квантовый на оси не знагленные целочисленные значения или токен или вложенный трюк с любым комбинацией рапольного тензора F4E2M1FN или типа F6E2M3FN или F6E3M2FN Тип или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FN или тип F8E4M3FNUZ или тип F8E4M3B11FNUZ или тип F8E5M2 или F8E5M2FNUZ или тип F8E8M0FNU или 16-битный плавучий или 32-битный плавучий или 64-битный плавучий или BFLOAT16-тип. -bit integer) или 2/4/8/16/32/64-битный неверный целый integer или 2/4/8/16/32/64-разрядный не знаковый целое число или сложный тип с 32-битным поплавковым или 64-битным поплавковым элементом или 2/4/8/16/32-битный равномерный квантовый квантовый целое число или 2/4/8/16/32-бита. на оси, подписанное целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаменитые значения целочисленных или значений токена

mhlo.concatenate (mhlo :: concatenateop)

Конкатенатная операция

Соглашает переменного числа тензоров в inputs вдоль dimension в том же порядке, что и заданные аргументы, и дает тензор result .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#concatenate

Пример:

%result = mhlo.concatenate %input0, %input1, dim = 0 : (tensor<3x2xi64>, tensor<1x2xi64>) -> tensor<4x2xi64>

Черты: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Интерфейсы: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
dimension :: mlir :: integerattr 64-разрядный знаменитый целочисленный атрибут, значение которого нетрицательное.

Операнды:

Операнд Описание
val Вариада по ранжированию типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или типа F8E3M4 или типа F8E4M3B11FNE8 16-битный поплавок или 32-битный поплавок или 64- битовая плавание или тип Bfloat16 или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-разрядочное неверное целое число или 2/4/8/16/32/64-разрядное целое число или комплекс Введите с 32-разрядными платными или 64-битными поплавковыми элементами или 2/4/8/16/32-битным равномерным квантовым подписанным целым числом или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8 /16/32-битный равномерный квантовый на оси.

Результаты:

Результат Описание
«Безымянный» Ранг Тензор типа F4E2M1FN или типа F6E2M3FN или типа F6E3M2FN или тип F8E3M4 или тип F8E4M3 или тип F8E4M3FNUZ или тип F8E4M3FNUZ 32-битный поплавок или 64-битный поплавок или bfloat16 тип или Pred (aka boolean или 1-битное целое число) или 2/4/8/16/32/64-битное целое число или 2/4/8/16/32/64-бита без знака или комплексный тип с 32-разрядный поплавок или 64-разрядные поплавковые элементы или 2/4/8/16/32-битное равномерное квантовое целое число или 2/4/8/16/32-битное равномерное квантовое целое число без знака или 2/4/8/16 /32-битная равномерная квантованая на оси целое число или 2/4/8/16/32-битное равномерное квантовое на оси не знаковые целочисленные значения

mhlo.constant (mhlo :: constantop)

Постоянная операция

Производит output тензор из постоянного value .

См.: Https://github.com/openxla/stablehlo/blob/main/docs/spec.md#constant

Пример:

%output = mhlo.constant dense<[[0.0, 1.0], [2.0, 3.0]]> : tensor<2x2xf32>

Черты: AlwaysSpeculatableImplTrait , ConstantLike

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Эффекты: MemoryEffects::Effect{}

Атрибуты:

Атрибут Тип Mlir Описание
value :: mlir :: elementsattr Постоянный атрибут вектора/тензора

Результаты:

Результат Описание
output statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.convert (mhlo::ConvertOp)

Convert operation

Синтаксис:

operation ::= `mhlo.convert` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs an element-wise conversion from one element type to another on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convert

Пример:

%result = mhlo.convert %operand : (tensor<3xi32>) -> tensor<3xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.convolution (mhlo::ConvolutionOp)

Convolution operation

Синтаксис:

operation ::= `mhlo.convolution` `(`operands`)`
              `dim_numbers` `=` custom<ConvolutionDimensions>($dimension_numbers) `,`
              `window` `=` `{` custom<WindowAttributes>($window_strides, $padding,
              $lhs_dilation, $rhs_dilation,
              $window_reversal) `}`
              attr-dict `:` functional-type(operands, results)

Computes dot products between windows of lhs and slices of rhs and produces result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convolution

Пример:

%result = "mhlo.convolution"(%lhs, %rhs) {
  window_strides = dense<4> : tensor<2xi64>,
  padding = dense<0> : tensor<2x2xi64>,
  lhs_dilation = dense<2> : tensor<2xi64>,
  rhs_dilation = dense<1> : tensor<2xi64>,
  window_reversal = dense<false> : tensor<2xi1>,
  dimension_numbers = #mhlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi32>, tensor<3x3x1x1xi32>) -> tensor<1x2x2x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
lhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
rhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_reversal ::mlir::DenseElementsAttr constant boolean vector/tensor attribute
dimension_numbers ::mlir::mhlo::ConvDimensionNumbersAttr Structure of dimension information for conv op
feature_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.copy (mhlo::CopyOp)

Copy operation

Синтаксис:

operation ::= `mhlo.copy` operands attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation a copy of operand . Depending on the metadata attached to the operation, it can behave quite differently from a no-op.

Пример:

%0 = mhlo.copy %arg0 : tensor<f32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
cross_program_prefetch_index ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.cosine (mhlo::CosineOp)

Cosine operation

Синтаксис:

operation ::= `mhlo.cosine` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise cosine operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cosine

Пример:

%result = mhlo.cosine %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.count_leading_zeros (mhlo::ClzOp)

Clz operation

Синтаксис:

operation ::= `mhlo.count_leading_zeros` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise count of the number of leading zero bits in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#count_leading_zeros

Пример:

%result = mhlo.count_leading_zeros %operand : tensor<2x2xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.create_token (mhlo::CreateTokenOp)

CreateToken operation

Синтаксис:

operation ::= `mhlo.create_token` attr-dict `:` type(results)

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as AfterAllOp with 0 inputs: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all

Пример:

%output = mhlo.create_token : !mhlo.token

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Результаты:

Результат Описание
output жетон

mhlo.cross-replica-sum (mhlo::CrossReplicaSumOp)

CrossReplicaSum operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as AllReduceOp with channel_id = 0 , use_global_device_ids = false and computation implementing addition: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce

Пример:

%result = "mhlo.cross-replica-sum"(%operand) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
} : (tensor<4xf32>) -> tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.custom_call (mhlo::CustomCallOp)

CustomCall operation

Синтаксис:

operation ::= `mhlo.custom_call` custom<CustomCallTarget>($call_target_name) `(` $inputs `)`
              attr-dict `:` functional-type(operands, results)

Encapsulates an implementation-defined operation call_target_name that takes inputs and called_computations and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#custom_call

Пример:

%results = "mhlo.custom_call"(%input0) {
  call_target_name = "foo",
  has_side_effect = false,
  backend_config = "bar",
  api_version = 1 : i32,
  called_computations = [@foo]
} : (tensor<f32>) -> tensor<f32>

A custom call invokes code external to XLA. The `inputs` are passed to the
external code, and the external code is expected to produce a result of the
given type. The exact mechanism is backend-specific. For example, in the CPU
backend, a call instruction is emitted which targets a symbol with the name
`call_target_name`.

If XLA runtime is enabled for a backend, then custom calls use the runtime
custom call calling convention to call into the external functions. This
calling convention defines an ABI for encoding arguments, attributes and
results.

Depending on the API version there are two ways to pass extra bits of static
information to the external function:

1. For `API_VERSION_TYPED_FFI` custom calls `backend_config` must be a
   dictionary attribute, that will be encoded according to the custom call
   calling convention and passed to the external function as the attributes
   argument. External code is expected to use declarative bindings (see
   `xla/runtime/custom_call.h`) to decode them at run time. These custom
   calls are only supported if XLA uses XLA runtime.

2. For previous API versions it is the user responsibility to encode extra
   bits of static information as a string `backend_config` attribute, and
   decode it at run time.

Interfaces: MemoryEffectOpInterface

Attributes:

Атрибут MLIR Type Описание
call_target_name ::mlir::StringAttr string attribute
has_side_effect ::mlir::BoolAttr bool attribute
backend_config ::mlir::Attribute string attribute or dictionary of named attribute values
api_version ::mlir::mhlo::CustomCallApiVersionAttr Custom call API version
called_computations ::mlir::ArrayAttr flat symbol ref array attribute
custom_call_schedule ::mlir::mhlo::CustomCallScheduleAttr Specifies the desired schedule for the custom-call.
operand_layouts ::mlir::ArrayAttr Array of layout (1D tensor of index type) attributes
result_layouts ::mlir::ArrayAttr Array of layout (1D tensor of index type) attributes
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of CustomCall

Operands:

Операнд Описание
inputs variadic of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
«unnamed» variadic of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.divide (mhlo::DivOp)

Div operation

Синтаксис:

operation ::= `mhlo.divide` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise division of dividend lhs and divisor rhs tensors and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#divide

Пример:

%result = mhlo.divide %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.domain (mhlo::DomainOp)

Domain operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, these operations are used to group instructions with the same DomainMetadata property. ShardingMetadata is the main use case today to group instructions on the same device. Domain instructions provide two major benefits:

  • Prevent unintentionally optimizing instructions across domains.
  • Automatically assign the metadata of the instructions created in the domain. Without domain instructions, each HLO optimization pass would have to check and propagate the metadata, which would be easy to miss and also adds complexity to the compiler. Since domain instructions connect two different domains, each domain instruction is associated with two DomainMetadata -- one on the operand side and one on the user side of the domain.

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
kind ::mlir::mhlo::DomainKindAttr Kind of domain metatdata attached to an HLO domain.
entry_metadata ::mlir::StringAttr string attribute
exit_metadata ::mlir::StringAttr string attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.dot (mhlo::DotOp)

Dot operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as XLA's Dot: https://www.tensorflow.org/xla/operation_semantics#dot

Пример:

%0 = mhlo.dot %arg0, %arg1 : (tensor<1x2xi32>, tensor<2x1xi32>) -> tensor<1x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dot_general (mhlo::DotGeneralOp)

DotGeneral operation

Computes dot products between slices of lhs and slices of rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dot_general

Пример:

%result = "mhlo.dot_general"(%lhs, %rhs) {
  dot_dimension_numbers = #mhlo.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>]
} : (tensor<2x2x2xi32>, tensor<2x2x2xi32>) -> tensor<2x2x2xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dot_dimension_numbers ::mlir::mhlo::DotDimensionNumbersAttr Attribute that models the dimension information for dot.
precision_config ::mlir::ArrayAttr Precision Config attribute
algorithm ::mlir::mhlo::DotAlgorithmAttr Attribute that models the algorithm constraints to use for computing dot.

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_broadcast_in_dim (mhlo::DynamicBroadcastInDimOp)

DynamicBroadcastInDim operation

This operation is functionally identical to broadcast_in_dim op, but the result shape is specified dynamically via output_dimensions .

It also accepts optional attributes to express static knowledge about the expanding behavior of dimensions. If not specified, all dimensions are assumed to be possibly expanding. The sets of dimensions that are known to be expanding and the set of dimensions that are known to be non-expanding must be disjoint and they must be a subset of the operand's dimensions.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_broadcast_in_dim

Пример:

%operand = mhlo.constant dense<[[1, 2, 3]]> : tensor<1x3xi64>
%output_dimensions = mhlo.constant dense<[2, 3, 2]> : tensor<3xi64>
%result = "mhlo.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>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
broadcast_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
known_expanding_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
known_nonexpanding_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
output_dimensions 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_conv (mhlo::DynamicConvOp)

DynamicConv operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as ConvolutionOp except that padding is specified dynamically via d_padding : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convolution

Пример:

%result = "mhlo.dynamic_conv"(%lhs, %rhs, %d_padding) {
  window_strides = dense<4> : tensor<2xi64>,
  lhs_dilation = dense<2> : tensor<2xi64>,
  rhs_dilation = dense<1> : tensor<2xi64>,
  window_reversal = dense<false> : tensor<2xi1>,
  dimension_numbers = #mhlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi32>, tensor<3x3x1x1xi32>, tensor<2x2xi64>) -> tensor<1x2x2x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
lhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
rhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_reversal ::mlir::DenseElementsAttr constant boolean vector/tensor attribute
dimension_numbers ::mlir::mhlo::ConvDimensionNumbersAttr Structure of dimension information for conv op
feature_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
d_padding ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_gather (mhlo::DynamicGatherOp)

DynamicGather operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as GatherOp except that slice_sizes are specified dynamically: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#gather

Пример:

%result = "mhlo.dynamic_gather"(%operand, %start_indices, %slice_sizes) {
  dimension_numbers = #mhlo.gather<
    offset_dims = [2, 3],
    collapsed_slice_dims = [0],
    start_index_map = [0, 2],
    index_vector_dim = 2>,
  indices_are_sorted = false
} : (tensor<3x4x2xi32>, tensor<2x3x2xi64>, tensor<3xi64>) -> tensor<2x3x2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
slice_sizes statically shaped 1-dimensional integer tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_iota (mhlo::DynamicIotaOp)

DynamicIota operation

This operation is functionally identical to iota op, but the result shape is specified dynamically via output_shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_iota

Пример:

%0 = mhlo.dynamic_iota %arg0, dim = 0 : (tensor<1xindex>) -> tensor<4xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
output_shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_pad (mhlo::DynamicPadOp)

DynamicPad operation

Синтаксис:

operation ::= `mhlo.dynamic_pad` operands attr-dict `:` functional-type(operands, results)

Dynamically Pads the operand , with amount of padding added at low-end/high-end/interior is passed through input tensors.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
edge_padding_low 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
edge_padding_high 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
interior_padding 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_reshape (mhlo::DynamicReshapeOp)

DynamicReshape operation

Синтаксис:

operation ::= `mhlo.dynamic_reshape` operands attr-dict `:` functional-type(operands, results)

This operation is functionally identical to reshape op, but the result shape is specified dynamically via output_shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_reshape

Пример:

%output_shape = mhlo.constant dense<[3, 2]> : tensor<2xi64>
%result = mhlo.dynamic_reshape %operand, %output_shape : (tensor<2x3xi64>, tensor<2xi64>) -> tensor<3x2xi64>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
output_shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_slice (mhlo::DynamicSliceOp)

DynamicSlice operation

Extracts a slice from the operand using dynamically-computed starting indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_slice

Пример:

%result = mhlo.dynamic_slice %operand, %start_indices0, %start_indices1, sizes = [2, 2]
  : (tensor<4x4xi32>, tensor<i64>, tensor<i64>) -> tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_update_slice (mhlo::DynamicUpdateSliceOp)

DynamicUpdateSlice operation

Синтаксис:

operation ::= `mhlo.dynamic_update_slice` operands attr-dict `:` functional-type(operands, results)

Produces a result tensor which is equal to the operand tensor except that the slice starting at start_indices is updated with the values in update .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_update_slice

Пример:

%result = mhlo.dynamic_update_slice %operand, %update, %start_indices0, %start_indices1
  : (tensor<4x4xi32>, tensor<2x2xi32>, tensor<i64>, tensor<i64>) -> tensor<4x4xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
update ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.einsum (mhlo::EinsumOp)

Einsum operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as TF's einsum: https://www.tensorflow.org/api_docs/python/tf/einsum

Пример:

%result = "mhlo.einsum"(%lhs, %rhs) {
  einsum_config = "ab,bc->ac"
} : (tensor<4x16xf32>, tensor<16x4xf32>) -> tensor<4x4xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
einsum_config ::mlir::StringAttr string attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.erf (mhlo::ErfOp)

Erf operation

Синтаксис:

operation ::= `mhlo.erf` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise erf operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#erf

Пример:

%result = mhlo.erf %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.exponential (mhlo::ExpOp)

Exp operation

Синтаксис:

operation ::= `mhlo.exponential` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise exponential operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#exponential

Пример:

%result = mhlo.exponential %operand : tensor<2x2xf64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.exponential_minus_one (mhlo::Expm1Op)

Expm1 operation

Синтаксис:

operation ::= `mhlo.exponential_minus_one` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise exponential minus one operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#exponential_minus_one

Пример:

%result = mhlo.exponential_minus_one %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.fft (mhlo::FftOp)

Fft operation

Performs the forward and inverse Fourier transforms for real and complex inputs/outputs.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#fft

Пример:

%result = mhlo.fft %operand, type = FFT, length = [4] : (tensor<4xcomplex<f32>>) -> tensor<4xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
fft_type ::mlir::mhlo::FftTypeAttr XLA fast fourier transform type.
fft_length ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.floor (mhlo::FloorOp)

Floor operation

Синтаксис:

operation ::= `mhlo.floor` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise floor of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#floor

Пример:

%result = mhlo.floor %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.fusion (mhlo::FusionOp)

Fusion operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation consists of a group of basic ops (represented as a region attached to it). It serves as a hint to the backend that it is beneficial to emit the contained ops into a single loop nest or kernel.

Attributes:

Атрибут MLIR Type Описание
fusion_kind ::mlir::mhlo::FusionKindAttr fusion kind
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of Fusion

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Результаты:

Результат Описание
results variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.gather (mhlo::GatherOp)

Gather operation

Gathers slices from operand tensor from offsets specified in start_indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#gather

Пример:

%result = "mhlo.gather"(%operand, %start_indices) {
  dimension_numbers = #stablehlo.gather<
    offset_dims = [3, 4],
    collapsed_slice_dims = [1],
    operand_batching_dims = [0],
    start_indices_batching_dims = [1],
    start_index_map = [2, 1],
    index_vector_dim = 3>,
  slice_sizes = dense<[0, 2, 2]> : tensor<3xi64>,
  indices_are_sorted = false
} : (tensor<2x3x4x2xi64>, tensor<2x2x3x2xi64>) -> tensor<2x2x3x2x2xi64>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.get_dimension_size (mhlo::GetDimensionSizeOp)

GetDimensionSize operation

Produces the size of the given dimension of the operand .

See https://github.com/openxla/stablehlo/blob/main/docs/spec.md#get_dimension_size

Пример:

%result = mhlo.get_dimension_size %operand, dim = 1 : (tensor<2x3xf32>) -> tensor<i32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» tensor of 32-bit signless integer values

mhlo.get_tuple_element (mhlo::GetTupleElementOp)

GetTupleElement operation

Синтаксис:

operation ::= `mhlo.get_tuple_element` $operand `[` $index `]` attr-dict `:` functional-type(operands, results)

Extracts element at index position of the operand tuple and produces a result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#get_tuple_element

Пример:

%result = mhlo.get_tuple_element %operand[0] : (tuple<tensor<2xf32>, tuple<tensor<i32>>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
index ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.if (mhlo::IfOp)

If operation

Produces the output from executing exactly one branch from true_branch or false_branch depending on the value of pred .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#if

Example: %result = "mhlo.if"(%pred) ({ "mhlo.return"(%result_true_branch) : (tensor ) -> () }, { "mhlo.return"(%result_false_branch) : (tensor ) -> () }) : (tensor ) -> tensor

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface

Operands:

Операнд Описание
pred ranked tensor of pred (AKA boolean or 1-bit integer) values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.imag (mhlo::ImagOp)

Imag operation

Синтаксис:

operation ::= `mhlo.imag` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Extracts the imaginary part, element-wise, from the operand and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#imag

Пример:

%result = mhlo.imag %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.infeed (mhlo::InfeedOp)

Infeed operation

Reads data from the infeed and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#infeed

Пример:

%results:2 = "mhlo.infeed"(%token) {
  infeed_config = ""
} : (!mhlo.token) -> (tensor<3x3x3xi32>, !mhlo.token)

Attributes:

Атрибут MLIR Type Описание
infeed_config ::mlir::StringAttr string attribute
layout ::mlir::ArrayAttr array attribute

Operands:

Операнд Описание
token жетон

Результаты:

Результат Описание
«unnamed» variadic of statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.iota (mhlo::IotaOp)

Iota operation

Fills an output tensor with values in increasing order starting from zero along the iota_dimension dimension.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#iota

Пример:

%output = mhlo.iota dim = 0 : tensor<4x5xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Результаты:

Результат Описание
output statically shaped tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements ценности

mhlo.is_finite (mhlo::IsFiniteOp)

IsFinite operation

Синтаксис:

operation ::= `mhlo.is_finite` $x attr-dict `:` functional-type(operands, results)

Performs element-wise check whether the value in x is finite (ie is neither +Inf, -Inf, nor NaN) and produces a y tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#is_finite

Пример:

%y = mhlo.is_finite %x : (tensor<7xf32>) -> tensor<7xi1>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
x ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
y ranked tensor of pred (AKA boolean or 1-bit integer) values

mhlo.log (mhlo::LogOp)

Log operation

Синтаксис:

operation ::= `mhlo.log` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logarithm operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#log

Пример:

%result = mhlo.log %operand : tensor<2x2xf64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.log_plus_one (mhlo::Log1pOp)

Log1p operation

Синтаксис:

operation ::= `mhlo.log_plus_one` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logarithm plus one operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#log_plus_one

Пример:

%result = mhlo.log_plus_one %operand : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.logistic (mhlo::LogisticOp)

Logistic operation

Синтаксис:

operation ::= `mhlo.logistic` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logistic operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#logistic

Пример:

%result = mhlo.logistic %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.map (mhlo::MapOp)

Map operation

Applies a map function computation to inputs along the dimensions and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#map

Пример:

%result = "mhlo.map"(%input0, %input1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.multiply %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  dimensions = dense<[0, 1]> : tensor<2xi64>
} : (tensor<2x2xi32>, tensor<2x2xi32>) -> tensor<2x2xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameOperandsAndResultShape , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.maximum (mhlo::MaxOp)

Max operation

Синтаксис:

operation ::= `mhlo.maximum` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise max operation on tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#maximum

Пример:

%result = mhlo.maximum %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.minimum (mhlo::MinOp)

Min operation

Синтаксис:

operation ::= `mhlo.minimum` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise min operation on tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#minimum

Пример:

%result = mhlo.minimum %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.minimum_broadcast_shapes (mhlo::MinimumBroadcastShapesOp)

Minimizes the rank of two or more shapes to be broadcasted

Синтаксис:

operation ::= `mhlo.minimum_broadcast_shapes` $shapes attr-dict `:` type($shapes) `->` type($results)

Given two or more 1D tensors representing shapes, returns one 1D tensor for each operand, where operand i corresponds to output i .

The returned tensors have the property that they specify a shape which is a reshape of the corresponding input shape, and the broadcasted output shape (using shape::BroadcastOp) of the returned shapes is a reshape of the broadcasted output shape of the input shapes. Among all possibilities with this property, the one is chosen which minimizes the rank of each returned shape.

The general idea of this op is that it can be used for ops which have a broadcasting semantic to operate on shapes with a possibly smaller rank while preserving equivalence of the computed values. After computing the result of the op using reshaped operands, the result can be reshaped to the result that would have been originally computed.

Here is an example with two input shapes:

mhlo.minimum_broadcast_shapes [1, 2, 3, 1, 2, 1],
                                 [1, 1, 1, 2, 3] -> [6, 2, 1], [2, 3]

The broadcasted output shape of the operands is [1, 2, 3, 1, 2, 3], the broadcasted output shape of the outputs is [6, 2, 3]. These two shapes are reshapes of each other, and also each output is a reshape of the corresponding input.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
shapes variadic of 1D tensor of index values

Результаты:

Результат Описание
results variadic of 1D tensor of index values

mhlo.multiply (mhlo::MulOp)

Mul operation

Синтаксис:

operation ::= `mhlo.multiply` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise product of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#multiply

Пример:

%result = mhlo.multiply %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.negate (mhlo::NegOp)

Neg operation

Синтаксис:

operation ::= `mhlo.negate` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise negation of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#negate

Пример:

%result = mhlo.negate %operand : tensor<2x3xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.not (mhlo::NotOp)

Not operation

Синтаксис:

operation ::= `mhlo.not` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise NOT of tensor operand of type integer and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#not

Пример:

%result = mhlo.not %operand : tensor<5x3x1xi1>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.optimization_barrier (mhlo::OptimizationBarrierOp)

OptimizationBarrier operation

Синтаксис:

operation ::= `mhlo.optimization_barrier` attr-dict ($operand^ `:` custom<PairwiseOpType>(type($operand), type($result))):(`(` `)`)?

Ensures that the operations that produce the operand are executed before any operations that depend on the result and prevents compiler transformations from moving operations across the barrier. Other than that, the operation is an identity, ie result = operand .

See https://github.com/openxla/stablehlo/blob/main/docs/spec.md#optimization_barrier

Пример:

%result0, %result1 = mhlo.optimization_barrier %operand0, %operand1 : tensor<f32>, tensor<f32>

Traits: AlwaysSpeculatableImplTrait , HLO_PairwiseSameOperandAndResultType

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Результаты:

Результат Описание
result variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.or (mhlo::OrOp)

Or operation

Синтаксис:

operation ::= `mhlo.or` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise OR of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#or

Пример:

%result = mhlo.or %lhs, %rhs : tensor<2xi1>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.outfeed (mhlo::OutfeedOp)

Outfeed operation

Writes inputs to the outfeed and produces a result token.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#outfeed

Пример:

%result = "mhlo.outfeed"(%input0, %token) {
  outfeed_config = ""
} : (tensor<3x3x3xi32>, !mhlo.token) -> !mhlo.token

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
outfeed_config ::mlir::StringAttr string attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
token жетон

Результаты:

Результат Описание
«unnamed» жетон

mhlo.pad (mhlo::PadOp)

Pad operation

Expands operand by padding around the tensor as well as between the elements of the tensor with the given padding_value .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#pad

Пример:

%0 = mhlo.pad %arg0, %arg1, low = [0, 1], high = [2, 1], interior = [1, 2]
  : (tensor<2x3xi32>, tensor<i32>) -> tensor<5x9xi32>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
edge_padding_low ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
edge_padding_high ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
interior_padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.partition_id (mhlo::PartitionIdOp)

PartitionId operation

Синтаксис:

operation ::= `mhlo.partition_id` attr-dict `:` type(results)

Produces partition_id of the current process.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#partition_id

Пример:

%result = mhlo.partition_id : tensor<ui32>

Interfaces: InferTypeOpInterface

Результаты:

Результат Описание
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.popcnt (mhlo::PopulationCountOp)

PopulationCount operation

Синтаксис:

operation ::= `mhlo.popcnt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise count of the number of bits set in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#popcnt

Пример:

%result = mhlo.popcnt %operand : tensor<4xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.power (mhlo::PowOp)

Pow operation

Синтаксис:

operation ::= `mhlo.power` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise exponentiation of lhs tensor by rhs tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#power

Пример:

%result = mhlo.power %lhs, %rhs : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.real (mhlo::RealOp)

Real operation

Синтаксис:

operation ::= `mhlo.real` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Extracts the real part, element-wise, from the operand and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#real

Пример:

%result = mhlo.real %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.real_dynamic_slice (mhlo::RealDynamicSliceOp)

RealDynamicSlice operation

Синтаксис:

operation ::= `mhlo.real_dynamic_slice` operands attr-dict `:` functional-type(operands, results)

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as SliceOp except that start_indices , limit_indices and strides are specified dynamically: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#slice

Пример:

%result = mhlo.real_dynamic_slice %operand,
            %start_indices, %limit_indices, %strides
       : (tensor<256x?xf32>, tensor<2xindex>, tensor<2xindex>, tensor<2xindex>) -> tensor<256x?xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
limit_indices 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
strides 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.recv (mhlo::RecvOp)

Recv operation

Receives data from a channel with channel_id and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#recv

Пример:

%results:2 = "mhlo.recv"(%token) {
  // channel_id = 5 : i64,
  // channel_type = #stablehlo<channel_type HOST_TO_DEVICE>,
  channel_handle = #mhlo.channel_handle<handle = 5, type = 3>,
  is_host_transfer = true
} : (!mhlo.token) -> (tensor<3x4xi32>, !mhlo.token)

Attributes:

Атрибут MLIR Type Описание
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
token жетон

Результаты:

Результат Описание
«unnamed» variadic of statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.reduce (mhlo::ReduceOp)

Reduce operation

Applies a reduction function body to inputs and init_values along the dimensions and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce

Пример:

%result = "mhlo.reduce"(%input, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0) : (tensor<i32>) -> ()
}) {
  dimensions = dense<1> : tensor<1xi64>
} : (tensor<1x6xi32>, tensor<i32>) -> tensor<1xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameVariadicOperandSize , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_values variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.reduce_precision (mhlo::ReducePrecisionOp)

ReducePrecision operation

Синтаксис:

operation ::= `mhlo.reduce_precision` $operand `,` `format` `=` custom<ExponentMantissa>($exponent_bits, $mantissa_bits)
              attr-dict `:` custom<SameOperandsAndResultType>(type($operand), type($output))

Performs element-wise conversion of operand to another floating-point type that uses exponent_bits and mantissa_bits and back to the original floating-point type and produces an output tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_precision

Пример:

%output = mhlo.reduce_precision %operand, format = e5m2 : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
exponent_bits ::mlir::IntegerAttr 32-bit signless integer attribute whose value is positive
mantissa_bits ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.reduce_scatter (mhlo::ReduceScatterOp)

ReduceScatter operation

Within each process group in the process grid, performs reduction, using computations , over the values of the operand tensor from each process, splits the reduction result along scatter_dimension into parts, and scatters the split parts between the processes to produce the result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_scatter

Пример:

%result = "mhlo.reduce_scatter"(%operand) ({
  ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>):
  %0 = mhlo.add %arg0, %arg1 : tensor<f32>
  mhlo.return %0 : tensor<f32>
}) {
  scatter_dimension = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<2x4xf32>) -> tensor<2x2xf32>

Attributes:

Атрибут MLIR Type Описание
scatter_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
use_global_device_ids ::mlir::UnitAttr unit attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.reduce_window (mhlo::ReduceWindowOp)

ReduceWindow operation

Applies a reduction function body to windows of inputs and init_values and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_window

Пример:

%result = "mhlo.reduce_window"(%input, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.add %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  window_dimensions = dense<[2, 1]> : tensor<2xi64>,
  window_strides = dense<[4, 1]> : tensor<2xi64>,
  base_dilations = dense<[2, 1]> : tensor<2xi64>,
  window_dilations = dense<[3, 1]> : tensor<2xi64>,
  padding = dense<[[2, 1], [0, 0]]> : tensor<2x2xi64>
} : (tensor<3x2xi32>, tensor<i32>) -> tensor<2x2xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameVariadicOperandSize , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
window_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
base_dilations ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_dilations ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_values variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.remainder (mhlo::RemOp)

Rem operation

Синтаксис:

operation ::= `mhlo.remainder` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise remainder of dividend lhs and divisor rhs tensors and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#remainder

Пример:

%result = mhlo.remainder %lhs, %rhs : tensor<4xi64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.replica_id (mhlo::ReplicaIdOp)

ReplicaId operation

Синтаксис:

operation ::= `mhlo.replica_id` attr-dict `:` type(results)

Produces replica_id of the current process.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#replica_id

Пример:

%result = mhlo.replica_id : tensor<ui32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Результаты:

Результат Описание
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.reshape (mhlo::ReshapeOp)

Reshape operation

Синтаксис:

operation ::= `mhlo.reshape` operands attr-dict `:` functional-type(operands, results)

Performs reshape of operand tensor to a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reshape

Пример:

%result = mhlo.reshape %operand : (tensor<2xf32>) -> tensor<1x2xf32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.return (mhlo::ReturnOp)

_This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/425

Informally, this operation serves as a terminator for regions defined by
the StableHLO ops. Non-StableHLO ops, e.g. `func.func`, have their own
terminators, e.g. `func.return`.

Example:

    ```mlir
    %result = "mhlo.reduce"(%input, %init_value) ({
      ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
        %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
        "mhlo.return"(%0) : (tensor<i32>) -> ()
    }) {
      dimensions = dense<1> : tensor<1xi64>
    } : (tensor<1x6xi32>, tensor<i32>) -> tensor<1xi32>
    ```_


Syntax:

```

operation ::= mhlo.return $results attr-dict ( : type($results)^)?



Traits: `AlwaysSpeculatableImplTrait`, `Terminator`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `results` | variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values


### `mhlo.reverse` (mhlo::ReverseOp)

_Reverse operation_

Reverses the order of elements in the `operand` along the specified
`dimensions` and produces a `result` tensor.

See:
<a href="https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reverse">https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reverse</a>

Example:
```mlir
%result = mhlo.reverse %operand, dims = [1] : tensor<3x2xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.rng (mhlo::RngOp)

Rng operation

Generates random numbers using the rng_distribution algorithm and produces a result tensor of a given shape shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rng

Пример:

%result = mhlo.rng %a, %b, %shape, distribution = NORMAL : (tensor<i32>, tensor<i32>, tensor<2xi64>) -> tensor<3x3xi32>

Traits: InferTensorType

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
rng_distribution ::mlir::mhlo::RngDistributionAttr XLA PRNG distribution to be used.

Operands:

Операнд Описание
a 0D tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
b 0D tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.rng_bit_generator (mhlo::RngBitGeneratorOp)

RngBitGenerator operation

Returns an output filled with uniform random data and an updated output state output_state given an initial state initial_state using the pseudorandom number generator algorithm rng_algorithm .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rng_bit_generator

Пример:

%output_state, %output = mhlo.rng_bit_generator %initial_state, algorithm = THREE_FRY : (tensor<2xui64>) -> (tensor<2xui64>, tensor<2x2xui64>)

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
rng_algorithm ::mlir::mhlo::RngAlgorithmAttr XLA PRNG algorithm to be used.

Operands:

Операнд Описание
initial_state ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
output_state ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
output statically shaped tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.round_nearest_afz (mhlo::RoundOp)

Round operation

Синтаксис:

operation ::= `mhlo.round_nearest_afz` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise rounding towards the nearest integer, breaking ties away from zero, on the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#round_nearest_afz

Пример:

%result = mhlo.round_nearest_afz %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.round_nearest_even (mhlo::RoundNearestEvenOp)

RoundNearestEven operation

Синтаксис:

operation ::= `mhlo.round_nearest_even` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise rounding towards the nearest integer, breaking ties towards the even integer, on the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#round_nearest_even

Пример:

%result = mhlo.round_nearest_even %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.rsqrt (mhlo::RsqrtOp)

Rsqrt operation

Синтаксис:

operation ::= `mhlo.rsqrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise reciprocal square root operation on operand tensor and produces a result tensor, implementing the rSqrt operation from the IEEE-754 specification.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rsqrt

Пример:

%result = mhlo.rsqrt %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.scatter (mhlo::ScatterOp)

Scatter operation

Produces results tensors which are equal to inputs tensors except that several slices specified by scatter_indices are updated with the values updates using update_computation .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#scatter

Пример:

%result = "mhlo.scatter"(%input, %scatter_indices, %update) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.add %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  scatter_dimension_numbers = #mhlo.scatter<
    update_window_dims = [3, 4],
    inserted_window_dims = [1],
    input_batching_dims = [0],
    scatter_indices_batching_dims = [1],
    scatter_dims_to_operand_dims = [2, 1],
    index_vector_dim = 3>,
  indices_are_sorted = false,
  unique_indices = false
} : (tensor<2x3x4x2xi64>, tensor<2x2x3x2xi64>, tensor<2x2x3x2x2xi64>) -> tensor<2x3x4x2xi64>

Traits: RecursiveMemoryEffects , SameVariadicOperandSize

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
scatter_dimension_numbers ::mlir::mhlo::ScatterDimensionNumbersAttr Attribute that models the dimension information for scatter
indices_are_sorted ::mlir::BoolAttr bool attribute
unique_indices ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
scatter_indices ranked tensor of integer or index values
updates variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.select (mhlo::SelectOp)

Select operation

Синтаксис:

operation ::= `mhlo.select` operands attr-dict `:`
              custom<SelectOpType>(type($pred), type($on_true), type($on_false), type($result))

Produces a result tensor where each element is selected from on_true or on_false tensor based on the value of the corresponding element of pred .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#select

Пример:

%result = mhlo.select %pred, %on_true, %on_false : tensor<2x2xi1>, tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_BroadcastingElementwise , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
pred ranked tensor of pred (AKA boolean or 1-bit integer) values
on_true ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
on_false ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.select_and_scatter (mhlo::SelectAndScatterOp)

SelectAndScatter operation

Scatters the values from the source tensor using scatter based on the outcome of reduce_window of the input tensor using select and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#select_and_scatter

Пример:

%result = "mhlo.select_and_scatter"(%operand, %source, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GE>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.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>

Traits: RecursiveMemoryEffects

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
window_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
source ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.send (mhlo::SendOp)

Send operation

Sends inputs to a channel channel_id and produces a result token.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#send

Пример:

%result = "mhlo.send"(%operand, %token) {
  // channel_id = 5 : i64,
  // channel_type = #stablehlo<channel_type DEVICE_TO_HOST>,
  channel_handle = #mhlo.channel_handle<handle = 5, type = 2>,
  is_host_transfer = true
} : (tensor<3x4xi32>, !mhlo.token) -> !mhlo.token

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
token жетон

Результаты:

Результат Описание
«unnamed» жетон

mhlo.set_dimension_size (mhlo::SetDimensionSizeOp)

SetDimensionSize operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as XLA's SetDimensionSize: https://www.tensorflow.org/xla/operation_semantics#setdimensionsize

Пример:

%0 = mhlo.set_dimension_size %arg0, %arg1, dim = 1 : (tensor<4x2xf32>, tensor<i32>) -> tensor<4x2xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
size tensor of 32-bit signless integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.shift_left (mhlo::ShiftLeftOp)

ShiftLeft operation

Синтаксис:

operation ::= `mhlo.shift_left` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise left-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_left

Пример:

%result = mhlo.shift_left %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.shift_right_arithmetic (mhlo::ShiftRightArithmeticOp)

ShiftRightArithmetic operation

Синтаксис:

operation ::= `mhlo.shift_right_arithmetic` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise arithmetic right-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_right_arithmetic

Пример:

%result = mhlo.shift_right_arithmetic %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.shift_right_logical (mhlo::ShiftRightLogicalOp)

ShiftRightLogical operation

Синтаксис:

operation ::= `mhlo.shift_right_logical` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise logical right-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_right_logical

Пример:

%result = mhlo.shift_right_logical %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.sign (mhlo::SignOp)

Sign operation

Синтаксис:

operation ::= `mhlo.sign` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Returns the sign of the operand element-wise and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sign

Пример:

%result = mhlo.sign %operand : tensor<7xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.sine (mhlo::SineOp)

Sine operation

Синтаксис:

operation ::= `mhlo.sine` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise sine operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sine

Пример:

%result = mhlo.sine %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.slice (mhlo::SliceOp)

Slice operation

Extracts a slice from the operand using statically-computed starting indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#slice

Пример:

%result = "mhlo.slice" (%operand) {
  start_indices = dense<[1, 2]> : tensor<2xi64>,
  limit_indices = dense<[3, 4]> : tensor<2xi64>,
  strides = dense<1> : tensor<2xi64>
} : (tensor<3x4xi64>) -> tensor<2x2xi64>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
start_indices ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
limit_indices ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sort (mhlo::SortOp)

Sort operation

Sorts a variadic number of tensors in inputs together, according to a custom comparator , along the given dimension and produces a variadic number of tensors as results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sort

Пример:

%result0, %result1 = "mhlo.sort"(%input0, %input1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>, %arg2: tensor<i32>, %arg3: tensor<i32>):
    %predicate = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GT>
      } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%predicate) : (tensor<i1>) -> ()
}) {
  dimension = 0 : i64,
  is_stable = true
} : (tensor<2x3xi32>, tensor<2x3xi32>) -> (tensor<2x3xi32>, tensor<2x3xi32>)

Traits: InferTensorType , RecursiveMemoryEffects , SameOperandsAndResultShape

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute
is_stable ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sparse_dot (mhlo::SparseDotOp)

Sparse dot operation

Similar to dot_general operation, with one or both of the operands being sparse. An additional argument provides sparsity meta information. Disclaimer: this op is experimental / a work in progress.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
lhs_sparsity ::mlir::mhlo::SparsityDescriptorAttr Describes structured (N:M) sparsity configuration
rhs_sparsity ::mlir::mhlo::SparsityDescriptorAttr Describes structured (N:M) sparsity configuration
dot_dimension_numbers ::mlir::mhlo::DotDimensionNumbersAttr Attribute that models the dimension information for dot.
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
meta variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sqrt (mhlo::SqrtOp)

Sqrt operation

Синтаксис:

operation ::= `mhlo.sqrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise square root operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sqrt

Пример:

%result = mhlo.sqrt %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.stochastic_convert (mhlo::StochasticConvertOp)

StochasticConvert operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/295

Informally, this operation performs element-wise conversion of values from a bigger type to a smaller one with stochastic rounding using the random number passed in.

Traits: AlwaysSpeculatableImplTrait , Elementwise

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
random ranked tensor of 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.subtract (mhlo::SubtractOp)

Subtract operation

Синтаксис:

operation ::= `mhlo.subtract` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise subtraction of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#subtract

Пример:

%result = mhlo.subtract %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.tan (mhlo::TanOp)

Tan operation

Синтаксис:

operation ::= `mhlo.tan` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/954

Informally, this operation returns Tan(operand) element-wise.

Пример:

%0 = mhlo.tan %arg0 : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.tanh (mhlo::TanhOp)

Tanh operation

Синтаксис:

operation ::= `mhlo.tanh` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise hyperbolic tangent operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#tanh

Пример:

%result = mhlo.tanh %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.topk (mhlo::TopKOp)

TopK operation

Синтаксис:

operation ::= `mhlo.topk` `(`$operand `,` `k` `=` $k (`,` `largest` `=` $largest^)? `)` attr-dict `:`
              type($operand) `->` `(`type($values)`,` type($indices)`)`

Returns top k values and their indices, along the last dimension of the operand if largest=true or the bottom k values if largest=false .

See: https://www.tensorflow.org/xla/operation_semantics#top-k

Пример:

%values, %indices = mhlo.topk(%operand, k=5, largest=true)
  : tensor<100xf32> -> (tensor<5xf32>, tensor<5xi32>)

Traits: InferTensorType , RecursiveMemoryEffects

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
k ::mlir::IntegerAttr 64-bit signless integer attribute
largest ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
values ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
indices ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.torch_index_select (mhlo::TorchIndexSelectOp)

TorchIndexSelect operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as PyTorch's index_select, augmented with support for batch dimensions: https://pytorch.org/docs/stable/generated/torch.index_select.html

The batch_dims attribute specifies the number of major batch dimensions (0 or more) that act like a multidimensional loop over both the operand and the index.

Пример:

%result = "mhlo.torch_index_select"(%operand, %index) {
  dim = 2 : i64,
  batch_dims = 1 : i64
} : (tensor<8x128x3072x64xf32>, tensor<8x16x1024xi32>) -> tensor<8x128x16x1024x64xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dim ::mlir::IntegerAttr 64-bit signless integer attribute
batch_dims ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
index ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.trace (mhlo::TraceOp)

Trace operation

Синтаксис:

operation ::= `mhlo.trace` $operand `,` $tag attr-dict `:` type($operand)

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/604

It is not used by JAX, PyTorch or TensorFlow, so it looks like we should've classified it as "Private to XLA" and not included it in StableHLO in the first place. With that in mind, its semantics will not be documented here.

Пример:

mhlo.trace %arg0, "In test code." : tensor<5x1x5xi32>

Attributes:

Атрибут MLIR Type Описание
tag ::mlir::StringAttr string attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.transpose (mhlo::TransposeOp)

Transpose operation

Permutes the dimensions of operand tensor using permutation and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#transpose

Пример:

%0 = mhlo.transpose %arg0, dims = [2, 1, 0] : (tensor<1x2x3xi32>) -> tensor<3x2x1xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
permutation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.triangular_solve (mhlo::TriangularSolveOp)

TriangularSolve operation

Solves batches of systems of linear equations with lower or upper triangular coefficient matrices.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#triangular_solve

Пример:

%result = "mhlo.triangular_solve"(%a, %b) {
  left_side = true,
  lower = true,
  unit_diagonal = false,
  transpose_a = #stablehlo<transpose NO_TRANSPOSE>
} : (tensor<3x3xf32>, tensor<3x3xf32>) -> tensor<3x3xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
left_side ::mlir::BoolAttr bool attribute
lower ::mlir::BoolAttr bool attribute
unit_diagonal ::mlir::BoolAttr bool attribute
transpose_a ::mlir::mhlo::TransposeAttr Transpose options

Operands:

Операнд Описание
a ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values
b ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.tuple (mhlo::TupleOp)

Tuple operation

Синтаксис:

operation ::= `mhlo.tuple` $val attr-dict `:` custom<TupleOpType>(type($val), type($result))

Produces a result tuple from values val .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#tuple

Пример:

%result = mhlo.tuple %val0, %val1 : tuple<tensor<2xf32>, tuple<tensor<i32>>>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
val variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
result nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.uniform_dequantize (mhlo::UniformDequantizeOp)

UniformDequantize operation

Синтаксис:

operation ::= `mhlo.uniform_dequantize` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise conversion of quantized tensor operand to a floating-point tensor result according to the quantization parameters defined by the operand type.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#uniform_dequantize

Пример:

%result = mhlo.uniform_dequantize %operand : (tensor<16x16x!quant.uniform<i8:f32, 34.0:16>>) -> tensor<16x16xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , InferTensorType , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.uniform_quantize (mhlo::UniformQuantizeOp)

UniformQuantize operation

Синтаксис:

operation ::= `mhlo.uniform_quantize` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise conversion of floating-point tensor or quantized tensor operand to a quantized tensor result according to the quantization parameters defined by the result type.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#uniform_quantize

Пример:

%result = mhlo.uniform_quantize %operand : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.while (mhlo::WhileOp)

While operation

Produces the output from executing body function 0 or more times while the cond function outputs true .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#while

Пример:

%results0, %results1 = "mhlo.while"(%operand0, %operand1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction LT>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %constant0) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0, %arg1) : (tensor<i32>, tensor<i32>) -> ()
}) : (tensor<i32>, tensor<i32>) -> (tensor<i32>, tensor<i32>)

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface , OpAsmOpInterface

Operands:

Операнд Описание
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.xla.rng_get_and_update_state (mhlo::XlaRngGetAndUpdateStateOp)

XlaRngGetAndUpdateState operation

Синтаксис:

operation ::= `mhlo.xla.rng_get_and_update_state` attr-dict

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation represents the change of the global random number generator state for rng instructions. The global state is incremented by delta and the old state is returned.

The output is currently defined for a single output type. If this changes in the future to support multiple types, lowering to use of a global memref must ensure that a single memref is still used and updated appropriately.

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
delta ::mlir::IntegerAttr 64-bit signless integer attribute

Результаты:

Результат Описание
«unnamed» statically shaped tensor of 64-bit unsigned integer values

mhlo.xor (mhlo::XorOp)

Xor operation

Синтаксис:

operation ::= `mhlo.xor` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise XOR of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#xor

Пример:

%result = mhlo.xor %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Атрибуты

ArgResultAliasAttr

Attribute that models the alias relationship of entry function argument

This attribute captures the alias relationship of an MHLO main function argument to one of the results, denoted by resultIndex . The argTupleIndices and resultTupleIndices are used to index into nested tuples in operand and result respectively. If isMustAlias is true then the operand-result pair must alias.

This is meant to be used as an attribute on a function argument in MHLO. For example, in the following code it expresses that %arg1 may alias 0-th result.

func @main(%arg0: tensor<2xf32>, %arg1: tensor<3xf32> {mhlo.result_alias =
    mhlo.result_alias<result_index = [2], ...>}
  ) -> tensor<2xf32>, tensor<3xf32> {
  // function body ...
}

Параметры:

Параметр C++ type Описание
argTupleIndices ::llvm::ArrayRef<int64_t> Измерение
resultIndex int64_t
resultTupleIndices ::llvm::ArrayRef<int64_t> Измерение
isMustAlias bool

ChannelHandleAttr

two 64-bit integers 'handle' and 'type'

Синтаксис:

#mhlo.channel_handle<
  int64_t,   # handle
  int64_t   # type
>

Параметры:

Параметр C++ type Описание
ручка int64_t
тип int64_t

ComparisonDirectionAttr

Which comparison operation to perform.

Синтаксис:

#mhlo.comparison_direction<
  ::mlir::mhlo::ComparisonDirection   # value
>

Enum cases:

  • EQ ( EQ )
  • NE ( NE )
  • GE ( GE )
  • GT ( GT )
  • LE ( LE )
  • LT ( LT )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::ComparisonDirection an enum of type ComparisonDirection

ComparisonTypeAttr

Which comparison type to use.

Синтаксис:

#mhlo.comparison_type<
  ::mlir::mhlo::ComparisonType   # value
>

Enum cases:

  • NOTYPE ( NOTYPE )
  • FLOAT ( FLOAT )
  • TOTALORDER ( TOTALORDER )
  • SIGNED ( SIGNED )
  • UNSIGNED ( UNSIGNED )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::ComparisonType an enum of type ComparisonType

ConvDimensionNumbersAttr

Structure of dimension information for conv op

Параметры:

Параметр C++ type Описание
inputBatchDimension int64_t
inputFeatureDimension int64_t
inputSpatialDimensions ::llvm::ArrayRef<int64_t> Измерение
kernelInputFeatureDimension int64_t
kernelOutputFeatureDimension int64_t
kernelSpatialDimensions ::llvm::ArrayRef<int64_t> Измерение
outputBatchDimension int64_t
outputFeatureDimension int64_t
outputSpatialDimensions ::llvm::ArrayRef<int64_t> Измерение

CrossProgramPrefetchAttr

Argument that is prefetched from another program

Синтаксис:

#mhlo.cross_program_prefetch<
  int64_t,   # parameter
  ::llvm::ArrayRef<int64_t>,   # indices
  std::optional<int64_t>   # offset
>

This attribute captures an argument that is prefetched from another program. For a given CrossProgramPrefetchAttr , parameter tells us which argument of the main function of the module is prefetched, and indices is a shape index telling us what subshape of that argument is prefetched.

A shape has a subshape iff it is a tuple. In that case, the subshape of the tuple by indices is the shape achieved after indexing by each element of indices in turn. For example, the [1,0] subshape of tuple<tuple<token, token>, tuple<tensor<i32>, token>> is tensor<i32> .

An empty value for indices means the whole shape is prefetched.

Например,

module attributes { mhlo.cross_program_prefetch = [ #mhlo.cross_program_prefetch< parameter = 0, indices = [0]> ]} {
  func.func @copy(%arg0 : tuple<tensor<2x3xi32>, tensor<i32>>) -> tuple<tensor<2x3xi32>, tensor<i32>> {
    %0 = "mhlo.copy"(%arg0) {is_cross_program_prefetch}
    return %0 : tuple<tensor<2x3xi32>, tensor<i32>>
  }
  func.func @main(%arg0 : tuple<tensor<2x3xi32>, tensor<i32>>) -> tuple<tensor<2x3xi32>, tensor<i32>> {
    %1 = "mhlo.async_start"(%arg0) {called_computation=@copy}
    %2 = "mhlo.async_done"(%1) {called_computation=@copy}
    return %2 : tuple<tensor<2x3xi32>, tensor<i32>>
  }
}

The parameter = 0 tells us that the async copy of the 0 th parameter is a cross_program_prefetch , while the index of [0] tells us that the 0 th element of the tuple is prefetched while the other element of the tuple is not.

Параметры:

Параметр C++ type Описание
параметр int64_t
индексы ::llvm::ArrayRef<int64_t> Измерение
компенсировать std::optional<int64_t>

CustomCallScheduleAttr

Specifies the desired schedule for the custom-call.

Синтаксис:

#mhlo.custom_call_schedule<
  ::mlir::mhlo::CustomCallSchedule   # value
>

Enum cases:

  • NONE ( NONE )
  • LATEST ( LATEST )
  • EARLIEST ( EARLIEST )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::CustomCallSchedule an enum of type CustomCallSchedule

DequantizeModeAttr

Dequantization mode. Only MIN_COMBINED is supported.

Синтаксис:

#mhlo.dequantize_mode<
  ::mlir::mhlo::DequantizeMode   # value
>

Enum cases:

  • MIN_COMBINED ( MIN_COMBINED )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::DequantizeMode an enum of type DequantizeMode

DomainKindAttr

Kind of domain metatdata attached to an HLO domain.

Синтаксис:

#mhlo.kind<
  ::mlir::mhlo::DomainKind   # value
>

Enum cases:

  • sharding ( sharding )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::DomainKind an enum of type DomainKind

DotAlgorithmAttr

Attribute that models the algorithm constraints to use for computing dot.

Синтаксис:

#mhlo.dot_algorithm<
  Type,   # lhsPrecisionType
  Type,   # rhsPrecisionType
  Type,   # accumulationType
  int64_t,   # lhsComponentCount
  int64_t,   # rhsComponentCount
  int64_t,   # numPrimitiveOperations
  bool   # allowImpreciseAccumulation
>

Параметры:

Параметр C++ type Описание
lhsPrecisionType Type
rhsPrecisionType Type
accumulationType Type
lhsComponentCount int64_t
rhsComponentCount int64_t
numPrimitiveOperations int64_t
allowImpreciseAccumulation bool

DotDimensionNumbersAttr

Attribute that models the dimension information for dot.

Параметры:

Параметр C++ type Описание
lhsBatchingDimensions ::llvm::ArrayRef<int64_t> Измерение
rhsBatchingDimensions ::llvm::ArrayRef<int64_t> Измерение
lhsContractingDimensions ::llvm::ArrayRef<int64_t> Измерение
rhsContractingDimensions ::llvm::ArrayRef<int64_t> Измерение

FftTypeAttr

XLA fast fourier transform type.

Синтаксис:

#mhlo.fft_type<
  ::mlir::mhlo::FftType   # value
>

Enum cases:

  • FFT ( FFT )
  • IFFT ( IFFT )
  • RFFT ( RFFT )
  • IRFFT ( IRFFT )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::FftType an enum of type FftType

FusionKindAttr

fusion kind

Синтаксис:

#mhlo.fusion_kind<
  ::mlir::mhlo::FusionKind   # value
>

Enum cases:

  • kLoop ( kLoop )
  • kInput ( kInput )
  • kOutput ( kOutput )
  • kCustom ( kCustom )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::FusionKind an enum of type FusionKind

GatherDimensionNumbersAttr

Attribute that models the dimension information for gather

Параметры:

Параметр C++ type Описание
offsetDims ::llvm::ArrayRef<int64_t> Измерение
collapsedSliceDims ::llvm::ArrayRef<int64_t> Измерение
operandBatchingDims ::llvm::ArrayRef<int64_t> Измерение
startIndicesBatchingDims ::llvm::ArrayRef<int64_t> Измерение
startIndexMap ::llvm::ArrayRef<int64_t> Измерение
indexVectorDim int64_t

OutputOperandAliasAttr

Attribute that models the alias relationship of output and operand of a CustomCall op

Синтаксис:

#mhlo.output_operand_alias<
  ::llvm::ArrayRef<int64_t>,   # outputTupleIndices
  int64_t,   # operandIndex
  ::llvm::ArrayRef<int64_t>   # operandTupleIndices
>

This attribute captures the alias relationship of the output to one of the operands for a CustomCall op, denoted by operand_index . The output_tuple_indices and operand_tuple_indices are used to index into output and operand types. These indices lists are empty if the corresponding types are not tuple types, and can be arbitrarily long in case of arbitrarily nested tuple types.

See https://www.tensorflow.org/xla/aliasing

Example when used as array with in mhlo.custom-call:

%0 = "mhlo.custom_call"(%arg0, %arg1) {
  // other attributes
  output_operand_alias = [
    #mhlo.output_operand_alias<output_tuple_indices = [0],
                               operand_index = 0,
                               operand_tuple_indices = [1]>
  ]
} : (tuple<tensor<1x1xf32>, tensor<2x3xf32>>, tensor<5x5xf32>) -> tuple<tensor<2x3xf32>>

The output and the 0th operand are both tuples. The aliasing shows the
relationship between the 0th element in output tuple with the 1st element in
the 0th operand. And both of them are of the same type: tensor<2x3xf32>.

Параметры:

Параметр C++ type Описание
outputTupleIndices ::llvm::ArrayRef<int64_t> Измерение
operandIndex int64_t
operandTupleIndices ::llvm::ArrayRef<int64_t> Измерение

PrecisionAttr

XLA precision for an operand. Has backend specific meaning.

Синтаксис:

#mhlo.precision<
  ::mlir::mhlo::Precision   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • HIGH ( HIGH )
  • HIGHEST ( HIGHEST )
  • PACKED_NIBBLE ( PACKED_NIBBLE )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::Precision an enum of type Precision

RngAlgorithmAttr

XLA PRNG algorithm to be used.

Синтаксис:

#mhlo.rng_algorithm<
  ::mlir::mhlo::RngAlgorithm   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • THREE_FRY ( THREE_FRY )
  • PHILOX ( PHILOX )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::RngAlgorithm an enum of type RngAlgorithm

RngDistributionAttr

XLA PRNG distribution to be used.

Синтаксис:

#mhlo.rng_distribution<
  ::mlir::mhlo::RngDistribution   # value
>

Enum cases:

  • UNIFORM ( UNIFORM )
  • NORMAL ( NORMAL )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::RngDistribution an enum of type RngDistribution

ScatterDimensionNumbersAttr

Attribute that models the dimension information for scatter

Параметры:

Параметр C++ type Описание
updateWindowDims ::llvm::ArrayRef<int64_t> Измерение
insertedWindowDims ::llvm::ArrayRef<int64_t> Измерение
inputBatchingDims ::llvm::ArrayRef<int64_t> Измерение
scatterIndicesBatchingDims ::llvm::ArrayRef<int64_t> Измерение
scatterDimsToOperandDims ::llvm::ArrayRef<int64_t> Измерение
indexVectorDim int64_t

SparsityDescriptorAttr

Describes structured (N:M) sparsity configuration

Синтаксис:

#mhlo.sparsity<
  int64_t,   # dimension
  int64_t,   # n
  int64_t   # m
>

This attribute is defined for a sparse dot operation with a structured sparse input tensor. With (N=2,M=4), every 4 consecutive logical elements have exactly 2 non-zero physical elements in the input tensor.

$dimension defines the index of the contracting dimension that is sparse (it has to be the most minor dimension). The additional metadata operand in the sparse dot operation defines which logical elements are zeroed out.

Параметры:

Параметр C++ type Описание
измерение int64_t
н int64_t
м int64_t

TransposeAttr

Transpose options

Синтаксис:

#mhlo.transpose<
  ::mlir::mhlo::Transpose   # value
>

Enum cases:

  • TRANSPOSE_INVALID ( TRANSPOSE_INVALID )
  • NO_TRANSPOSE ( NO_TRANSPOSE )
  • TRANSPOSE ( TRANSPOSE )
  • ADJOINT ( ADJOINT )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::Transpose an enum of type Transpose

TypeExtensionsAttr

Attribute that extends tensor type with MHLO type properties.

Синтаксис:

#mhlo.type_extensions<
  ::llvm::ArrayRef<int64_t>   # bounds
>

This attribute is used to extend MLIR tensor type with MHLO tensor specific properties. These properties aren't modeled in the MLIR type. This attribute is set in the encoding field of the tensor type.

See HLO_BoundedAttrInterface for documentation for bounds .

Параметры:

Параметр C++ type Описание
границы ::llvm::ArrayRef<int64_t>

Типы

AsyncBundleType

Opaque collection of other types

Синтаксис:

!mhlo.async_bundle<
  ::llvm::ArrayRef<Type>   # types
>

Параметры:

Параметр C++ type Описание
типы ::llvm::ArrayRef<Type>

Перечисления

ComparisonDirection

Which comparison operation to perform.

Случаи:

Символ Ценить Нить
эквалайзер 0 эквалайзер
СВ 1 СВ
GE 2 GE
ГТ 3 ГТ
ЛЕ 4 ЛЕ
LT 5 LT

ComparisonType

Which comparison type to use.

Случаи:

Символ Ценить Нить
NOTYPE 0 NOTYPE
ПЛАВАТЬ 1 ПЛАВАТЬ
TOTALORDER 2 TOTALORDER
ПОДПИСАНО 3 ПОДПИСАНО
UNSIGNED 4 UNSIGNED

CustomCallApiVersion

Custom call API version

Случаи:

Символ Ценить Нить
API_VERSION_UNSPECIFIED 0 API_VERSION_UNSPECIFIED
API_VERSION_ORIGINAL 1 API_VERSION_ORIGINAL
API_VERSION_STATUS_RETURNING 2 API_VERSION_STATUS_RETURNING
API_VERSION_STATUS_RETURNING_UNIFIED 3 API_VERSION_STATUS_RETURNING_UNIFIED
API_VERSION_TYPED_FFI 4 API_VERSION_TYPED_FFI

CustomCallSchedule

Specifies the desired schedule for the custom-call.

Случаи:

Символ Ценить Нить
НИКТО 0 НИКТО
ПОСЛЕДНИЙ 1 ПОСЛЕДНИЙ
EARLIEST 2 EARLIEST

DequantizeMode

Dequantization mode. Only MIN_COMBINED is supported.

Случаи:

Символ Ценить Нить
MIN_COMBINED 0 MIN_COMBINED

DomainKind

Kind of domain metatdata attached to an HLO domain.

Случаи:

Символ Ценить Нить
шардинг 0 шардинг

FftType

XLA fast fourier transform type.

Случаи:

Символ Ценить Нить
БПФ 0 БПФ
IFFT 1 IFFT
RFFT 2 RFFT
IRFFT 3 IRFFT

FusionKind

fusion kind

Случаи:

Символ Ценить Нить
kLoop 0 kLoop
kInput 1 kInput
kOutput 2 kOutput
kCustom 3 kCustom

Точность

XLA precision for an operand. Has backend specific meaning.

Случаи:

Символ Ценить Нить
ПО УМОЛЧАНИЮ 0 ПО УМОЛЧАНИЮ
ВЫСОКИЙ 1 ВЫСОКИЙ
HIGHEST 2 HIGHEST
PACKED_NIBBLE 3 PACKED_NIBBLE

RngAlgorithm

XLA PRNG algorithm to be used.

Случаи:

Символ Ценить Нить
ПО УМОЛЧАНИЮ 0 ПО УМОЛЧАНИЮ
THREE_FRY 1 THREE_FRY
PHILOX 2 PHILOX

RngDistribution

XLA PRNG distribution to be used.

Случаи:

Символ Ценить Нить
UNIFORM 1 UNIFORM
НОРМАЛЬНЫЙ 2 НОРМАЛЬНЫЙ

Transpose

Transpose options

Случаи:

Символ Ценить Нить
TRANSPOSE_INVALID 0 TRANSPOSE_INVALID
NO_TRANSPOSE 1 NO_TRANSPOSE
TRANSPOSE 2 TRANSPOSE
ADJOINT 3 ADJOINT
,

Операции

mhlo.abs (mhlo::AbsOp)

Abs operation

Синтаксис:

operation ::= `mhlo.abs` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise abs operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#abs

Пример:

%result = mhlo.abs %operand : tensor<3xi32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.add (mhlo::AddOp)

Add operation

Синтаксис:

operation ::= `mhlo.add` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise addition of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#add

Пример:

%result = mhlo.add %lhs, %rhs : tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.add_dependency (mhlo::AddDependencyOp)

AddDependency operation

Синтаксис:

operation ::= `mhlo.add_dependency` operands attr-dict `:` functional-type(operands, results)

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation two operands: a data operand and a token. The output of the operation is the data operand. When used with AfterAll this operation enables ordering non-side-effecting operations (those that do not produce token values).

Пример:

%1 = mhlo.add_dependency %arg0, %0 : (tensor<3x4xf32>, !mhlo.token) -> tensor<3x4xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token
token жетон

Результаты:

Результат Описание
output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.after_all (mhlo::AfterAllOp)

AfterAll operation

Синтаксис:

operation ::= `mhlo.after_all` $inputs attr-dict
              `:` custom<VariadicSameOperandsAndResultType>(ref($inputs), type($inputs), type($result))

Ensures that the operations producing the inputs are executed before any operations that depend on result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all

Пример:

%result = mhlo.after_all %input0, %input1 : !mhlo.token

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
inputs variadic of token

Результаты:

Результат Описание
result жетон

mhlo.all_gather (mhlo::AllGatherOp)

AllGather operation

Within each process group in the process grid, concatenates the values of the operand tensor from each process along all_gather_dim and produces a result tensor. The computation is applied separately for each operand in operands , producing one result per operand.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_gather

Пример:

%result = "mhlo.all_gather"(%operand) {
  all_gather_dim = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>,
  // use_global_device_ids = false
} : (tensor<2x2xf32>) -> tensor<2x4xf32>

Traits: SameOperandsAndResultElementType

Attributes:

Атрибут MLIR Type Описание
all_gather_dim ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
use_global_device_ids ::mlir::UnitAttr unit attribute

Operands:

Операнд Описание
operands variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.all_reduce (mhlo::AllReduceOp)

AllReduce operation

Within each process group in the process grid, applies a reduction function computation to the values of an operand tensor from each process and produces a result tensor. The computation is applied separately for each operand in operands , producing one result per operand.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce

Пример:

%result = "mhlo.all_reduce"(%operand) ({
  ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>):
    %0 = mhlo.add %arg1, %arg2 : tensor<f32>
    mhlo.return %0 : tensor<f32>
}) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<4xf32>) -> tensor<4xf32>

Traits: InferTensorType , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
use_global_device_ids ::mlir::UnitAttr unit attribute

Operands:

Операнд Описание
operands variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.all_to_all (mhlo::AllToAllOp)

AllToAll operation

Within each process group in the process grid, splits the values of the operand tensor along split_dimension into parts, scatters the split parts between the processes, concatenates the scattered parts along concat_dimension and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_to_all

Пример:

%result = "mhlo.all_to_all"(%operand) {
  split_dimension = 1 : i64,
  concat_dimension = 0 : i64,
  split_count = 2 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
} : (tensor<2x4xf32>) -> tensor<4x2xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsElementType , SameOperandsShape , SameVariadicOperandSize

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
split_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative
concat_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative
split_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'

Operands:

Операнд Описание
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.and (mhlo::AndOp)

And operation

Синтаксис:

operation ::= `mhlo.and` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise AND of two tensors lhs and rhs and produces a result tensor

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#and

Пример:

%result = mhlo.and %lhs, %rhs : tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.async_done (mhlo::AsyncDoneOp)

AsyncDone operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation blocks until the end of an asynchronous computation. It returns the final result of the asynchronous computation.

See the documentation for AsyncStart for more information.

Interfaces: InferTypeOpInterface

Operands:

Операнд Описание
bundle async_bundle with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.async_start (mhlo::AsyncStartOp)

AsyncStart operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation kicks off an asynchronous computation.

This is used when there are functions that contain both asynchronous waits (such as DMAs) and on-thread computation. For example, a function might consist of a computation, a DMA, another computation, a second DMA, and a final computation. This would be represented as an async_start followed by and async_update and an async_done. The async_start would do the first computation on-thread and then start the DMA. The async_update would wait for the DMA to complete if it wasn't yet done, then execute the second computation in the function, and start the second DMA. Finally, the async_done would wait on this last DMA, and then run the last computation that needs to be run on-thread and return the result of that final computation.

operands are passed to the computation directly called_computation is the function that will be run asynchronously execution_thread is the name of the thread in which it will be run. The main thread is called "main". All threads have names.

This returns all the state needed between async ops. After buffer assignment, the return values represents the space needed to hold the input, results, and any scratchpads needed or edited by the async op.

Attributes:

Атрибут MLIR Type Описание
called_computation ::mlir::FlatSymbolRefAttr flat symbol reference attribute
execution_thread ::mlir::StringAttr string attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
«unnamed» async_bundle with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.async_update (mhlo::AsyncUpdateOp)

AsyncUpdate operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation blocks on an asynchronous computation until a sync barrier. This returns bundle after operating on it.

See the documentation for AsyncStart for more information.

Interfaces: InferTypeOpInterface

Operands:

Операнд Описание
bundle async_bundle with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Результаты:

Результат Описание
«unnamed» async_bundle with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.atan2 (mhlo::Atan2Op)

Atan2 operation

Синтаксис:

operation ::= `mhlo.atan2` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise atan2 operation on lhs and rhs tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#atan2

Пример:

%result = mhlo.atan2 %lhs, %rhs : tensor<3xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.batch_norm_grad (mhlo::BatchNormGradOp)

BatchNormGrad operation

Computes gradients of several inputs of BatchNormTrainingOp backpropagating from grad_output , and produces grad_operand , grad_scale and grad_offset tensors.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_grad

Пример:

%grad_operand, %grad_scale, %grad_offset =
"mhlo.batch_norm_grad"(%operand, %scale, %mean, %variance, %grad_output) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>,
    tensor<2x2x2xf32>) -> (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>)

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
epsilon ::mlir::FloatAttr 32-bit float attribute
feature_index ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
scale 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
mean 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
variance 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
grad_output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
grad_operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
grad_scale 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
grad_offset 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.batch_norm_inference (mhlo::BatchNormInferenceOp)

BatchNormInference operation

Normalizes the operand tensor across all dimensions except for the feature_index dimension and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_inference

Пример:

%result = "mhlo.batch_norm_inference"(%operand, %scale, %offset, %mean, %variance) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>, tensor<2xf32>) -> tensor<2x2x2xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
epsilon ::mlir::FloatAttr 32-bit float attribute
feature_index ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
scale 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
offset 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
mean 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
variance 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.batch_norm_training (mhlo::BatchNormTrainingOp)

BatchNormTraining operation

Computes mean and variance across batch and spatial dimensions and normalizes the operand tensor, for each feature in the feature_index dimension and produces output , batch_mean and batch_var tensors.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#batch_norm_training

Пример:

%output, %batch_mean, %batch_var = "mhlo.batch_norm_training"(%operand, %scale, %offset) {
  epsilon = 0.0 : f32,
  feature_index = 2 : i64
} : (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>) -> (tensor<2x2x2xf32>, tensor<2xf32>, tensor<2xf32>)

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
epsilon ::mlir::FloatAttr 32-bit float attribute
feature_index ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
scale 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
offset 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Результаты:

Результат Описание
output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
batch_mean 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
batch_var 1D tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.bitcast (mhlo::BitcastOp)

Bitcast operation

Синтаксис:

operation ::= `mhlo.bitcast` operands attr-dict `:` functional-type(operands, results)

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation changes the shape of the input in the way that the physical arrangement of elements are unchanged.

This operation needs layout information to make sense of "physical arrangement of elements", and layout support in MHLO is currently a work in progress.

Пример:

%0 = mhlo.bitcast %arg0 : (tensor<3x4xf32>) -> tensor<3x4x1xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.bitcast_convert (mhlo::BitcastConvertOp)

BitcastConvert operation

Синтаксис:

operation ::= `mhlo.bitcast_convert` operands attr-dict `:` functional-type(operands, results)

Performs a bitcast operation on operand tensor and produces a result tensor where the bits of the entire operand tensor are reinterpreted using the type of the result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#bitcast_convert

Пример:

%result = mhlo.bitcast_convert %operand : (tensor<2xf32>) -> tensor<2x4xi8>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.broadcast (mhlo::BroadcastOp)

Операция вещания

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as XLA's Broadcast: https://www.tensorflow.org/xla/operation_semantics#broadcast

Пример:

%result = mhlo.broadcast %operand, sizes = [1, 2] : (tensor<3xi32>) -> tensor<1x2x3xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
broadcast_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.broadcast_in_dim (mhlo::BroadcastInDimOp)

BroadcastInDim operation

Expands the dimensions and/or rank of an input tensor by duplicating the data in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#broadcast_in_dim

Пример:

%result = mhlo.broadcast_in_dim %operand, dims = [2, 1] : (tensor<1x3xi32>) -> tensor<2x3x2xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
broadcast_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.case (mhlo::CaseOp)

Case operation

Produces the output from executing exactly one function from branches depending on the value of index .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#case

Пример:

%result0, %result1 = "mhlo.case"(%index) ({
  mhlo.return %result_branch0, %result_branch0 : tensor<2xi64>, tensor<2xi64>
}, {
  mhlo.return %result_branch1, %result_branch1 : tensor<2xi64>, tensor<2xi64>
}) : (tensor<i32>) -> (tensor<2xi64>, tensor<2xi64>)

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface

Operands:

Операнд Описание
index tensor of 32-bit signless integer values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.cbrt (mhlo::CbrtOp)

Cbrt operation

Синтаксис:

operation ::= `mhlo.cbrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise cubic root operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cbrt

Пример:

%result = mhlo.cbrt %operand : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.ceil (mhlo::CeilOp)

Ceil operation

Синтаксис:

operation ::= `mhlo.ceil` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise ceil of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#ceil

Пример:

%result = mhlo.ceil %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.cholesky (mhlo::CholeskyOp)

Cholesky operation

Computes the Cholesky decomposition of a batch of matrices.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cholesky

Пример:

%result = mhlo.cholesky %a, lower = true : tensor<3x3xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
lower ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
a ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.clamp (mhlo::ClampOp)

Clamp operation

Синтаксис:

operation ::= `mhlo.clamp` $min `,` $operand `,` $max attr-dict
              `:` custom<SameOperandsAndResultType>(type($min), type($operand), type($max), type($result))

Clamps every element of the operand tensor between a minimum and maximum value and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#clamp

Пример:

%result = mhlo.clamp %min, %operand, %max : tensor<3xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_BroadcastingElementwise , InferTensorType , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
min ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
max ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.collective_broadcast (mhlo::CollectiveBroadcastOp)

CollectiveBroadcast operation

Within each process group in the process grid, send the value of the operand tensor from the source process to the target processes and produce a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_broadcast

Пример:

%result = "mhlo.collective_broadcast"(%operand) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
} : (tensor<1x2xi64>) -> tensor<1x2xi64>

Traits: CompatibleOperandsAndResultType

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.collective_permute (mhlo::CollectivePermuteOp)

CollectivePermute operation

Within each process group in the process grid, sends the value of the operand tensor from the source process to the target process and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_permute

Пример:

%result = "mhlo.collective_permute"(%operand) {
  source_target_pairs = dense<[[0, 1], [1, 2]]> : tensor<2x2xi64>,
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
} : (tensor<4x2xf32>) -> tensor<4x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
source_target_pairs ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.compare (mhlo::CompareOp)

Compare operation

Синтаксис:

operation ::= `mhlo.compare` $comparison_direction `,` $lhs `,` $rhs (`,` $compare_type^)?
              attr-dict `:` functional-type(operands, results)

Performs element-wise comparison of lhs and rhs tensors according to comparison_direction and compare_type , and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#compare

Пример:

%result = mhlo.compare LT, %lhs, %rhs, FLOAT : (tensor<2xf32>, tensor<2xf32>) -> tensor<2xi1>

Traits: AlwaysSpeculatableImplTrait , Elementwise , InferTensorType , SameOperandsAndResultShape , SameOperandsElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
comparison_direction ::mlir::mhlo::ComparisonDirectionAttr Which comparison operation to perform.
compare_type ::mlir::mhlo::ComparisonTypeAttr Which comparison type to use.

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of pred (AKA boolean or 1-bit integer) values

mhlo.complex (mhlo::ComplexOp)

Сложная операция

Синтаксис:

operation ::= `mhlo.complex` operands attr-dict
              `:` custom<ComplexOpType>(type($lhs), type($rhs), type($result))

Performs element-wise conversion to a complex value from a pair of real and imaginary values, lhs and rhs , and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#complex

Пример:

%result = mhlo.complex %lhs, %rhs : tensor<2xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape , SameOperandsElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 32-bit float or 64-bit float values
rhs ranked tensor of 32-bit float or 64-bit float values

Results:

Результат Описание
result ranked tensor of complex type with 32-bit float or 64-bit float elements values

mhlo.composite (mhlo::CompositeOp)

Composite operation

Синтаксис:

operation ::= `mhlo.composite` $name $inputs attr-dict `:` functional-type(operands, results)

Encapsulates an operation made up (composed) of other StableHLO operations, taking inputs and composite_attributes and producing results . The semantics of the op are implemented by the decomposition attribute. The composite op can be replaced with its decomposition without changing program semantics. In cases where inlining the decomposition does not provide the same op semantics, prefer using custom_call .

The version field (defaults to 0 ) is used to denote when a composite's semantics change.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#composite

Пример:

%results = mhlo.composite "my.op" %arg0, %arg1 {
  decomposition = @my_op,
  composite_attributes = { my_attribute = "my_value" },
  version = 1 : i32
} : (tensor<f32>, tensor<f32>) -> tensor<f32>

Interfaces: SymbolUserOpInterface

Attributes:

Атрибут MLIR Type Описание
name ::mlir::StringAttr string attribute
composite_attributes ::mlir::DictionaryAttr dictionary of named attribute values
decomposition ::mlir::FlatSymbolRefAttr flat symbol reference attribute
version ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Results:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.concatenate (mhlo::ConcatenateOp)

Concatenate operation

Concatenates a variadic number of tensors in inputs along dimension dimension in the same order as the given arguments and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#concatenate

Пример:

%result = mhlo.concatenate %input0, %input1, dim = 0 : (tensor<3x2xi64>, tensor<1x2xi64>) -> tensor<4x2xi64>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
val variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.constant (mhlo::ConstantOp)

Constant operation

Produces an output tensor from a constant value .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#constant

Пример:

%output = mhlo.constant dense<[[0.0, 1.0], [2.0, 3.0]]> : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , ConstantLike

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
value ::mlir::ElementsAttr constant vector/tensor attribute

Результаты:

Результат Описание
output statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.convert (mhlo::ConvertOp)

Convert operation

Синтаксис:

operation ::= `mhlo.convert` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs an element-wise conversion from one element type to another on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convert

Пример:

%result = mhlo.convert %operand : (tensor<3xi32>) -> tensor<3xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.convolution (mhlo::ConvolutionOp)

Convolution operation

Синтаксис:

operation ::= `mhlo.convolution` `(`operands`)`
              `dim_numbers` `=` custom<ConvolutionDimensions>($dimension_numbers) `,`
              `window` `=` `{` custom<WindowAttributes>($window_strides, $padding,
              $lhs_dilation, $rhs_dilation,
              $window_reversal) `}`
              attr-dict `:` functional-type(operands, results)

Computes dot products between windows of lhs and slices of rhs and produces result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convolution

Пример:

%result = "mhlo.convolution"(%lhs, %rhs) {
  window_strides = dense<4> : tensor<2xi64>,
  padding = dense<0> : tensor<2x2xi64>,
  lhs_dilation = dense<2> : tensor<2xi64>,
  rhs_dilation = dense<1> : tensor<2xi64>,
  window_reversal = dense<false> : tensor<2xi1>,
  dimension_numbers = #mhlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi32>, tensor<3x3x1x1xi32>) -> tensor<1x2x2x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
lhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
rhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_reversal ::mlir::DenseElementsAttr constant boolean vector/tensor attribute
dimension_numbers ::mlir::mhlo::ConvDimensionNumbersAttr Structure of dimension information for conv op
feature_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.copy (mhlo::CopyOp)

Copy operation

Синтаксис:

operation ::= `mhlo.copy` operands attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation a copy of operand . Depending on the metadata attached to the operation, it can behave quite differently from a no-op.

Пример:

%0 = mhlo.copy %arg0 : tensor<f32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
cross_program_prefetch_index ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.cosine (mhlo::CosineOp)

Cosine operation

Синтаксис:

operation ::= `mhlo.cosine` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise cosine operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cosine

Пример:

%result = mhlo.cosine %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.count_leading_zeros (mhlo::ClzOp)

Clz operation

Синтаксис:

operation ::= `mhlo.count_leading_zeros` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise count of the number of leading zero bits in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#count_leading_zeros

Пример:

%result = mhlo.count_leading_zeros %operand : tensor<2x2xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.create_token (mhlo::CreateTokenOp)

CreateToken operation

Синтаксис:

operation ::= `mhlo.create_token` attr-dict `:` type(results)

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as AfterAllOp with 0 inputs: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#after_all

Пример:

%output = mhlo.create_token : !mhlo.token

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:

Результат Описание
output жетон

mhlo.cross-replica-sum (mhlo::CrossReplicaSumOp)

CrossReplicaSum operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as AllReduceOp with channel_id = 0 , use_global_device_ids = false and computation implementing addition: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#all_reduce

Пример:

%result = "mhlo.cross-replica-sum"(%operand) {
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>
} : (tensor<4xf32>) -> tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.custom_call (mhlo::CustomCallOp)

CustomCall operation

Синтаксис:

operation ::= `mhlo.custom_call` custom<CustomCallTarget>($call_target_name) `(` $inputs `)`
              attr-dict `:` functional-type(operands, results)

Encapsulates an implementation-defined operation call_target_name that takes inputs and called_computations and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#custom_call

Пример:

%results = "mhlo.custom_call"(%input0) {
  call_target_name = "foo",
  has_side_effect = false,
  backend_config = "bar",
  api_version = 1 : i32,
  called_computations = [@foo]
} : (tensor<f32>) -> tensor<f32>

A custom call invokes code external to XLA. The `inputs` are passed to the
external code, and the external code is expected to produce a result of the
given type. The exact mechanism is backend-specific. For example, in the CPU
backend, a call instruction is emitted which targets a symbol with the name
`call_target_name`.

If XLA runtime is enabled for a backend, then custom calls use the runtime
custom call calling convention to call into the external functions. This
calling convention defines an ABI for encoding arguments, attributes and
results.

Depending on the API version there are two ways to pass extra bits of static
information to the external function:

1. For `API_VERSION_TYPED_FFI` custom calls `backend_config` must be a
   dictionary attribute, that will be encoded according to the custom call
   calling convention and passed to the external function as the attributes
   argument. External code is expected to use declarative bindings (see
   `xla/runtime/custom_call.h`) to decode them at run time. These custom
   calls are only supported if XLA uses XLA runtime.

2. For previous API versions it is the user responsibility to encode extra
   bits of static information as a string `backend_config` attribute, and
   decode it at run time.

Interfaces: MemoryEffectOpInterface

Attributes:

Атрибут MLIR Type Описание
call_target_name ::mlir::StringAttr string attribute
has_side_effect ::mlir::BoolAttr bool attribute
backend_config ::mlir::Attribute string attribute or dictionary of named attribute values
api_version ::mlir::mhlo::CustomCallApiVersionAttr Custom call API version
called_computations ::mlir::ArrayAttr flat symbol ref array attribute
custom_call_schedule ::mlir::mhlo::CustomCallScheduleAttr Specifies the desired schedule for the custom-call.
operand_layouts ::mlir::ArrayAttr Array of layout (1D tensor of index type) attributes
result_layouts ::mlir::ArrayAttr Array of layout (1D tensor of index type) attributes
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of CustomCall

Operands:

Операнд Описание
inputs variadic of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Results:

Результат Описание
«unnamed» variadic of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.divide (mhlo::DivOp)

Div operation

Синтаксис:

operation ::= `mhlo.divide` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise division of dividend lhs and divisor rhs tensors and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#divide

Пример:

%result = mhlo.divide %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.domain (mhlo::DomainOp)

Domain operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, these operations are used to group instructions with the same DomainMetadata property. ShardingMetadata is the main use case today to group instructions on the same device. Domain instructions provide two major benefits:

  • Prevent unintentionally optimizing instructions across domains.
  • Automatically assign the metadata of the instructions created in the domain. Without domain instructions, each HLO optimization pass would have to check and propagate the metadata, which would be easy to miss and also adds complexity to the compiler. Since domain instructions connect two different domains, each domain instruction is associated with two DomainMetadata -- one on the operand side and one on the user side of the domain.

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
kind ::mlir::mhlo::DomainKindAttr Kind of domain metatdata attached to an HLO domain.
entry_metadata ::mlir::StringAttr string attribute
exit_metadata ::mlir::StringAttr string attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.dot (mhlo::DotOp)

Dot operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as XLA's Dot: https://www.tensorflow.org/xla/operation_semantics#dot

Пример:

%0 = mhlo.dot %arg0, %arg1 : (tensor<1x2xi32>, tensor<2x1xi32>) -> tensor<1x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dot_general (mhlo::DotGeneralOp)

DotGeneral operation

Computes dot products between slices of lhs and slices of rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dot_general

Пример:

%result = "mhlo.dot_general"(%lhs, %rhs) {
  dot_dimension_numbers = #mhlo.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>]
} : (tensor<2x2x2xi32>, tensor<2x2x2xi32>) -> tensor<2x2x2xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dot_dimension_numbers ::mlir::mhlo::DotDimensionNumbersAttr Attribute that models the dimension information for dot.
precision_config ::mlir::ArrayAttr Precision Config attribute
algorithm ::mlir::mhlo::DotAlgorithmAttr Attribute that models the algorithm constraints to use for computing dot.

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_broadcast_in_dim (mhlo::DynamicBroadcastInDimOp)

DynamicBroadcastInDim operation

This operation is functionally identical to broadcast_in_dim op, but the result shape is specified dynamically via output_dimensions .

It also accepts optional attributes to express static knowledge about the expanding behavior of dimensions. If not specified, all dimensions are assumed to be possibly expanding. The sets of dimensions that are known to be expanding and the set of dimensions that are known to be non-expanding must be disjoint and they must be a subset of the operand's dimensions.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_broadcast_in_dim

Пример:

%operand = mhlo.constant dense<[[1, 2, 3]]> : tensor<1x3xi64>
%output_dimensions = mhlo.constant dense<[2, 3, 2]> : tensor<3xi64>
%result = "mhlo.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>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
broadcast_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
known_expanding_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
known_nonexpanding_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
output_dimensions 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_conv (mhlo::DynamicConvOp)

DynamicConv operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as ConvolutionOp except that padding is specified dynamically via d_padding : https://github.com/openxla/stablehlo/blob/main/docs/spec.md#convolution

Пример:

%result = "mhlo.dynamic_conv"(%lhs, %rhs, %d_padding) {
  window_strides = dense<4> : tensor<2xi64>,
  lhs_dilation = dense<2> : tensor<2xi64>,
  rhs_dilation = dense<1> : tensor<2xi64>,
  window_reversal = dense<false> : tensor<2xi1>,
  dimension_numbers = #mhlo.conv<[b, 0, 1, f]x[0, 1, i, o]->[b, 0, 1, f]>,
  feature_group_count = 1 : i64,
  batch_group_count = 1 : i64,
  precision_config = [#stablehlo<precision DEFAULT>, #stablehlo<precision DEFAULT>]
} : (tensor<1x4x4x1xi32>, tensor<3x3x1x1xi32>, tensor<2x2xi64>) -> tensor<1x2x2x1xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
lhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
rhs_dilation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_reversal ::mlir::DenseElementsAttr constant boolean vector/tensor attribute
dimension_numbers ::mlir::mhlo::ConvDimensionNumbersAttr Structure of dimension information for conv op
feature_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
batch_group_count ::mlir::IntegerAttr 64-bit signless integer attribute whose value is positive
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
d_padding ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_gather (mhlo::DynamicGatherOp)

DynamicGather operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as GatherOp except that slice_sizes are specified dynamically: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#gather

Пример:

%result = "mhlo.dynamic_gather"(%operand, %start_indices, %slice_sizes) {
  dimension_numbers = #mhlo.gather<
    offset_dims = [2, 3],
    collapsed_slice_dims = [0],
    start_index_map = [0, 2],
    index_vector_dim = 2>,
  indices_are_sorted = false
} : (tensor<3x4x2xi32>, tensor<2x3x2xi64>, tensor<3xi64>) -> tensor<2x3x2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
slice_sizes statically shaped 1-dimensional integer tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_iota (mhlo::DynamicIotaOp)

DynamicIota operation

This operation is functionally identical to iota op, but the result shape is specified dynamically via output_shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_iota

Пример:

%0 = mhlo.dynamic_iota %arg0, dim = 0 : (tensor<1xindex>) -> tensor<4xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
output_shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_pad (mhlo::DynamicPadOp)

DynamicPad operation

Синтаксис:

operation ::= `mhlo.dynamic_pad` operands attr-dict `:` functional-type(operands, results)

Dynamically Pads the operand , with amount of padding added at low-end/high-end/interior is passed through input tensors.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
edge_padding_low 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
edge_padding_high 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
interior_padding 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_reshape (mhlo::DynamicReshapeOp)

DynamicReshape operation

Синтаксис:

operation ::= `mhlo.dynamic_reshape` operands attr-dict `:` functional-type(operands, results)

This operation is functionally identical to reshape op, but the result shape is specified dynamically via output_shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_reshape

Пример:

%output_shape = mhlo.constant dense<[3, 2]> : tensor<2xi64>
%result = mhlo.dynamic_reshape %operand, %output_shape : (tensor<2x3xi64>, tensor<2xi64>) -> tensor<3x2xi64>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
output_shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_slice (mhlo::DynamicSliceOp)

DynamicSlice operation

Extracts a slice from the operand using dynamically-computed starting indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_slice

Пример:

%result = mhlo.dynamic_slice %operand, %start_indices0, %start_indices1, sizes = [2, 2]
  : (tensor<4x4xi32>, tensor<i64>, tensor<i64>) -> tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.dynamic_update_slice (mhlo::DynamicUpdateSliceOp)

DynamicUpdateSlice operation

Синтаксис:

operation ::= `mhlo.dynamic_update_slice` operands attr-dict `:` functional-type(operands, results)

Produces a result tensor which is equal to the operand tensor except that the slice starting at start_indices is updated with the values in update .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#dynamic_update_slice

Пример:

%result = mhlo.dynamic_update_slice %operand, %update, %start_indices0, %start_indices1
  : (tensor<4x4xi32>, tensor<2x2xi32>, tensor<i64>, tensor<i64>) -> tensor<4x4xi32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
update ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices variadic of 0D tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.einsum (mhlo::EinsumOp)

Einsum operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as TF's einsum: https://www.tensorflow.org/api_docs/python/tf/einsum

Пример:

%result = "mhlo.einsum"(%lhs, %rhs) {
  einsum_config = "ab,bc->ac"
} : (tensor<4x16xf32>, tensor<16x4xf32>) -> tensor<4x4xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
einsum_config ::mlir::StringAttr string attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.erf (mhlo::ErfOp)

Erf operation

Синтаксис:

operation ::= `mhlo.erf` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise erf operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#erf

Пример:

%result = mhlo.erf %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.exponential (mhlo::ExpOp)

Exp operation

Синтаксис:

operation ::= `mhlo.exponential` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise exponential operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#exponential

Пример:

%result = mhlo.exponential %operand : tensor<2x2xf64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.exponential_minus_one (mhlo::Expm1Op)

Expm1 operation

Синтаксис:

operation ::= `mhlo.exponential_minus_one` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise exponential minus one operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#exponential_minus_one

Пример:

%result = mhlo.exponential_minus_one %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.fft (mhlo::FftOp)

Fft operation

Performs the forward and inverse Fourier transforms for real and complex inputs/outputs.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#fft

Пример:

%result = mhlo.fft %operand, type = FFT, length = [4] : (tensor<4xcomplex<f32>>) -> tensor<4xcomplex<f32>>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
fft_type ::mlir::mhlo::FftTypeAttr XLA fast fourier transform type.
fft_length ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.floor (mhlo::FloorOp)

Floor operation

Синтаксис:

operation ::= `mhlo.floor` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise floor of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#floor

Пример:

%result = mhlo.floor %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.fusion (mhlo::FusionOp)

Fusion operation

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation consists of a group of basic ops (represented as a region attached to it). It serves as a hint to the backend that it is beneficial to emit the contained ops into a single loop nest or kernel.

Attributes:

Атрибут MLIR Type Описание
fusion_kind ::mlir::mhlo::FusionKindAttr fusion kind
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of Fusion

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Results:

Результат Описание
results variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.gather (mhlo::GatherOp)

Gather operation

Gathers slices from operand tensor from offsets specified in start_indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#gather

Пример:

%result = "mhlo.gather"(%operand, %start_indices) {
  dimension_numbers = #stablehlo.gather<
    offset_dims = [3, 4],
    collapsed_slice_dims = [1],
    operand_batching_dims = [0],
    start_indices_batching_dims = [1],
    start_index_map = [2, 1],
    index_vector_dim = 3>,
  slice_sizes = dense<[0, 2, 2]> : tensor<3xi64>,
  indices_are_sorted = false
} : (tensor<2x3x4x2xi64>, tensor<2x2x3x2xi64>) -> tensor<2x2x3x2x2xi64>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.get_dimension_size (mhlo::GetDimensionSizeOp)

GetDimensionSize operation

Produces the size of the given dimension of the operand .

See https://github.com/openxla/stablehlo/blob/main/docs/spec.md#get_dimension_size

Пример:

%result = mhlo.get_dimension_size %operand, dim = 1 : (tensor<2x3xf32>) -> tensor<i32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» tensor of 32-bit signless integer values

mhlo.get_tuple_element (mhlo::GetTupleElementOp)

GetTupleElement operation

Синтаксис:

operation ::= `mhlo.get_tuple_element` $operand `[` $index `]` attr-dict `:` functional-type(operands, results)

Extracts element at index position of the operand tuple and produces a result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#get_tuple_element

Пример:

%result = mhlo.get_tuple_element %operand[0] : (tuple<tensor<2xf32>, tuple<tensor<i32>>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
index ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.if (mhlo::IfOp)

If operation

Produces the output from executing exactly one branch from true_branch or false_branch depending on the value of pred .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#if

Example: %result = "mhlo.if"(%pred) ({ "mhlo.return"(%result_true_branch) : (tensor ) -> () }, { "mhlo.return"(%result_false_branch) : (tensor ) -> () }) : (tensor ) -> tensor

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface

Operands:

Операнд Описание
pred ranked tensor of pred (AKA boolean or 1-bit integer) values

Результаты:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.imag (mhlo::ImagOp)

Imag operation

Синтаксис:

operation ::= `mhlo.imag` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Extracts the imaginary part, element-wise, from the operand and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#imag

Пример:

%result = mhlo.imag %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.infeed (mhlo::InfeedOp)

Infeed operation

Reads data from the infeed and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#infeed

Пример:

%results:2 = "mhlo.infeed"(%token) {
  infeed_config = ""
} : (!mhlo.token) -> (tensor<3x3x3xi32>, !mhlo.token)

Attributes:

Атрибут MLIR Type Описание
infeed_config ::mlir::StringAttr string attribute
layout ::mlir::ArrayAttr array attribute

Operands:

Операнд Описание
token жетон

Results:

Результат Описание
«unnamed» variadic of statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.iota (mhlo::IotaOp)

Iota operation

Fills an output tensor with values in increasing order starting from zero along the iota_dimension dimension.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#iota

Пример:

%output = mhlo.iota dim = 0 : tensor<4x5xi32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Results:

Результат Описание
output statically shaped tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements ценности

mhlo.is_finite (mhlo::IsFiniteOp)

IsFinite operation

Синтаксис:

operation ::= `mhlo.is_finite` $x attr-dict `:` functional-type(operands, results)

Performs element-wise check whether the value in x is finite (ie is neither +Inf, -Inf, nor NaN) and produces a y tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#is_finite

Пример:

%y = mhlo.is_finite %x : (tensor<7xf32>) -> tensor<7xi1>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
x ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Results:

Результат Описание
y ranked tensor of pred (AKA boolean or 1-bit integer) values

mhlo.log (mhlo::LogOp)

Log operation

Синтаксис:

operation ::= `mhlo.log` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logarithm operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#log

Пример:

%result = mhlo.log %operand : tensor<2x2xf64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.log_plus_one (mhlo::Log1pOp)

Log1p operation

Синтаксис:

operation ::= `mhlo.log_plus_one` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logarithm plus one operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#log_plus_one

Пример:

%result = mhlo.log_plus_one %operand : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.logistic (mhlo::LogisticOp)

Logistic operation

Синтаксис:

operation ::= `mhlo.logistic` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise logistic operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#logistic

Пример:

%result = mhlo.logistic %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.map (mhlo::MapOp)

Map operation

Applies a map function computation to inputs along the dimensions and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#map

Пример:

%result = "mhlo.map"(%input0, %input1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.multiply %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  dimensions = dense<[0, 1]> : tensor<2xi64>
} : (tensor<2x2xi32>, tensor<2x2xi32>) -> tensor<2x2xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameOperandsAndResultShape , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.maximum (mhlo::MaxOp)

Max operation

Синтаксис:

operation ::= `mhlo.maximum` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise max operation on tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#maximum

Пример:

%result = mhlo.maximum %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.minimum (mhlo::MinOp)

Min operation

Синтаксис:

operation ::= `mhlo.minimum` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise min operation on tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#minimum

Пример:

%result = mhlo.minimum %lhs, %rhs : tensor<4xf32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.minimum_broadcast_shapes (mhlo::MinimumBroadcastShapesOp)

Minimizes the rank of two or more shapes to be broadcasted

Синтаксис:

operation ::= `mhlo.minimum_broadcast_shapes` $shapes attr-dict `:` type($shapes) `->` type($results)

Given two or more 1D tensors representing shapes, returns one 1D tensor for each operand, where operand i corresponds to output i .

The returned tensors have the property that they specify a shape which is a reshape of the corresponding input shape, and the broadcasted output shape (using shape::BroadcastOp) of the returned shapes is a reshape of the broadcasted output shape of the input shapes. Among all possibilities with this property, the one is chosen which minimizes the rank of each returned shape.

The general idea of this op is that it can be used for ops which have a broadcasting semantic to operate on shapes with a possibly smaller rank while preserving equivalence of the computed values. After computing the result of the op using reshaped operands, the result can be reshaped to the result that would have been originally computed.

Here is an example with two input shapes:

mhlo.minimum_broadcast_shapes [1, 2, 3, 1, 2, 1],
                                 [1, 1, 1, 2, 3] -> [6, 2, 1], [2, 3]

The broadcasted output shape of the operands is [1, 2, 3, 1, 2, 3], the broadcasted output shape of the outputs is [6, 2, 3]. These two shapes are reshapes of each other, and also each output is a reshape of the corresponding input.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
shapes variadic of 1D tensor of index values

Results:

Результат Описание
results variadic of 1D tensor of index values

mhlo.multiply (mhlo::MulOp)

Mul operation

Синтаксис:

operation ::= `mhlo.multiply` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise product of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#multiply

Пример:

%result = mhlo.multiply %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.negate (mhlo::NegOp)

Neg operation

Синтаксис:

operation ::= `mhlo.negate` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise negation of operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#negate

Пример:

%result = mhlo.negate %operand : tensor<2x3xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.not (mhlo::NotOp)

Not operation

Синтаксис:

operation ::= `mhlo.not` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise NOT of tensor operand of type integer and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#not

Пример:

%result = mhlo.not %operand : tensor<5x3x1xi1>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.optimization_barrier (mhlo::OptimizationBarrierOp)

OptimizationBarrier operation

Синтаксис:

operation ::= `mhlo.optimization_barrier` attr-dict ($operand^ `:` custom<PairwiseOpType>(type($operand), type($result))):(`(` `)`)?

Ensures that the operations that produce the operand are executed before any operations that depend on the result and prevents compiler transformations from moving operations across the barrier. Other than that, the operation is an identity, ie result = operand .

See https://github.com/openxla/stablehlo/blob/main/docs/spec.md#optimization_barrier

Пример:

%result0, %result1 = mhlo.optimization_barrier %operand0, %operand1 : tensor<f32>, tensor<f32>

Traits: AlwaysSpeculatableImplTrait , HLO_PairwiseSameOperandAndResultType

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Results:

Результат Описание
result variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.or (mhlo::OrOp)

Or operation

Синтаксис:

operation ::= `mhlo.or` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise OR of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#or

Пример:

%result = mhlo.or %lhs, %rhs : tensor<2xi1>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.outfeed (mhlo::OutfeedOp)

Outfeed operation

Writes inputs to the outfeed and produces a result token.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#outfeed

Пример:

%result = "mhlo.outfeed"(%input0, %token) {
  outfeed_config = ""
} : (tensor<3x3x3xi32>, !mhlo.token) -> !mhlo.token

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
outfeed_config ::mlir::StringAttr string attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
token жетон

Results:

Результат Описание
«unnamed» жетон

mhlo.pad (mhlo::PadOp)

Pad operation

Expands operand by padding around the tensor as well as between the elements of the tensor with the given padding_value .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#pad

Пример:

%0 = mhlo.pad %arg0, %arg1, low = [0, 1], high = [2, 1], interior = [1, 2]
  : (tensor<2x3xi32>, tensor<i32>) -> tensor<5x9xi32>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
edge_padding_low ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
edge_padding_high ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
interior_padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
padding_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.partition_id (mhlo::PartitionIdOp)

PartitionId operation

Синтаксис:

operation ::= `mhlo.partition_id` attr-dict `:` type(results)

Produces partition_id of the current process.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#partition_id

Пример:

%result = mhlo.partition_id : tensor<ui32>

Interfaces: InferTypeOpInterface

Results:

Результат Описание
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.popcnt (mhlo::PopulationCountOp)

PopulationCount operation

Синтаксис:

operation ::= `mhlo.popcnt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise count of the number of bits set in the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#popcnt

Пример:

%result = mhlo.popcnt %operand : tensor<4xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.power (mhlo::PowOp)

Pow operation

Синтаксис:

operation ::= `mhlo.power` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise exponentiation of lhs tensor by rhs tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#power

Пример:

%result = mhlo.power %lhs, %rhs : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.real (mhlo::RealOp)

Real operation

Синтаксис:

operation ::= `mhlo.real` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Extracts the real part, element-wise, from the operand and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#real

Пример:

%result = mhlo.real %operand : (tensor<2xcomplex<f32>>) -> tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.real_dynamic_slice (mhlo::RealDynamicSliceOp)

RealDynamicSlice operation

Синтаксис:

operation ::= `mhlo.real_dynamic_slice` operands attr-dict `:` functional-type(operands, results)

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as SliceOp except that start_indices , limit_indices and strides are specified dynamically: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#slice

Пример:

%result = mhlo.real_dynamic_slice %operand,
            %start_indices, %limit_indices, %strides
       : (tensor<256x?xf32>, tensor<2xindex>, tensor<2xindex>, tensor<2xindex>) -> tensor<256x?xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
start_indices 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
limit_indices 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
strides 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.recv (mhlo::RecvOp)

Recv operation

Receives data from a channel with channel_id and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#recv

Пример:

%results:2 = "mhlo.recv"(%token) {
  // channel_id = 5 : i64,
  // channel_type = #stablehlo<channel_type HOST_TO_DEVICE>,
  channel_handle = #mhlo.channel_handle<handle = 5, type = 3>,
  is_host_transfer = true
} : (!mhlo.token) -> (tensor<3x4xi32>, !mhlo.token)

Attributes:

Атрибут MLIR Type Описание
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
token жетон

Results:

Результат Описание
«unnamed» variadic of statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.reduce (mhlo::ReduceOp)

Reduce operation

Applies a reduction function body to inputs and init_values along the dimensions and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce

Пример:

%result = "mhlo.reduce"(%input, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0) : (tensor<i32>) -> ()
}) {
  dimensions = dense<1> : tensor<1xi64>
} : (tensor<1x6xi32>, tensor<i32>) -> tensor<1xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameVariadicOperandSize , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_values variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.reduce_precision (mhlo::ReducePrecisionOp)

ReducePrecision operation

Синтаксис:

operation ::= `mhlo.reduce_precision` $operand `,` `format` `=` custom<ExponentMantissa>($exponent_bits, $mantissa_bits)
              attr-dict `:` custom<SameOperandsAndResultType>(type($operand), type($output))

Performs element-wise conversion of operand to another floating-point type that uses exponent_bits and mantissa_bits and back to the original floating-point type and produces an output tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_precision

Пример:

%output = mhlo.reduce_precision %operand, format = e5m2 : tensor<6xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
exponent_bits ::mlir::IntegerAttr 32-bit signless integer attribute whose value is positive
mantissa_bits ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Results:

Результат Описание
output ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.reduce_scatter (mhlo::ReduceScatterOp)

ReduceScatter operation

Within each process group in the process grid, performs reduction, using computations , over the values of the operand tensor from each process, splits the reduction result along scatter_dimension into parts, and scatters the split parts between the processes to produce the result .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_scatter

Пример:

%result = "mhlo.reduce_scatter"(%operand) ({
  ^bb0(%arg0: tensor<f32>, %arg1: tensor<f32>):
  %0 = mhlo.add %arg0, %arg1 : tensor<f32>
  mhlo.return %0 : tensor<f32>
}) {
  scatter_dimension = 1 : i64,
  replica_groups = dense<[[0, 1]]> : tensor<1x2xi64>,
  // channel_id = 0
  channel_handle = #mhlo.channel_handle<handle = 0, type = 0>
  // use_global_device_ids = false
} : (tensor<2x4xf32>) -> tensor<2x2xf32>

Attributes:

Атрибут MLIR Type Описание
scatter_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
use_global_device_ids ::mlir::UnitAttr unit attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.reduce_window (mhlo::ReduceWindowOp)

ReduceWindow operation

Applies a reduction function body to windows of inputs and init_values and produces results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reduce_window

Пример:

%result = "mhlo.reduce_window"(%input, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.add %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  window_dimensions = dense<[2, 1]> : tensor<2xi64>,
  window_strides = dense<[4, 1]> : tensor<2xi64>,
  base_dilations = dense<[2, 1]> : tensor<2xi64>,
  window_dilations = dense<[3, 1]> : tensor<2xi64>,
  padding = dense<[[2, 1], [0, 0]]> : tensor<2x2xi64>
} : (tensor<3x2xi32>, tensor<i32>) -> tensor<2x2xi32>

Traits: InferTensorType , RecursiveMemoryEffects , SameVariadicOperandSize , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
window_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
base_dilations ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_dilations ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_values variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.remainder (mhlo::RemOp)

Rem operation

Синтаксис:

operation ::= `mhlo.remainder` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise remainder of dividend lhs and divisor rhs tensors and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#remainder

Пример:

%result = mhlo.remainder %lhs, %rhs : tensor<4xi64>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.replica_id (mhlo::ReplicaIdOp)

ReplicaId operation

Синтаксис:

operation ::= `mhlo.replica_id` attr-dict `:` type(results)

Produces replica_id of the current process.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#replica_id

Пример:

%result = mhlo.replica_id : tensor<ui32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:

Результат Описание
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.reshape (mhlo::ReshapeOp)

Reshape operation

Синтаксис:

operation ::= `mhlo.reshape` operands attr-dict `:` functional-type(operands, results)

Performs reshape of operand tensor to a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reshape

Пример:

%result = mhlo.reshape %operand : (tensor<2xf32>) -> tensor<1x2xf32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» statically shaped tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.return (mhlo::ReturnOp)

_This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/425

Informally, this operation serves as a terminator for regions defined by
the StableHLO ops. Non-StableHLO ops, e.g. `func.func`, have their own
terminators, e.g. `func.return`.

Example:

    ```mlir
    %result = "mhlo.reduce"(%input, %init_value) ({
      ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
        %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
        "mhlo.return"(%0) : (tensor<i32>) -> ()
    }) {
      dimensions = dense<1> : tensor<1xi64>
    } : (tensor<1x6xi32>, tensor<i32>) -> tensor<1xi32>
    ```_


Syntax:

```

operation ::= mhlo.return $results attr-dict ( : type($results)^)?



Traits: `AlwaysSpeculatableImplTrait`, `Terminator`

Interfaces: `ConditionallySpeculatable`, `NoMemoryEffect (MemoryEffectOpInterface)`

Effects: `MemoryEffects::Effect{}`

#### Operands:

| Operand | Description |
| :-----: | ----------- |
| `results` | variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values


### `mhlo.reverse` (mhlo::ReverseOp)

_Reverse operation_

Reverses the order of elements in the `operand` along the specified
`dimensions` and produces a `result` tensor.

See:
<a href="https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reverse">https://github.com/openxla/stablehlo/blob/main/docs/spec.md#reverse</a>

Example:
```mlir
%result = mhlo.reverse %operand, dims = [1] : tensor<3x2xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.rng (mhlo::RngOp)

Rng operation

Generates random numbers using the rng_distribution algorithm and produces a result tensor of a given shape shape .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rng

Пример:

%result = mhlo.rng %a, %b, %shape, distribution = NORMAL : (tensor<i32>, tensor<i32>, tensor<2xi64>) -> tensor<3x3xi32>

Traits: InferTensorType

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
rng_distribution ::mlir::mhlo::RngDistributionAttr XLA PRNG distribution to be used.

Operands:

Операнд Описание
a 0D tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
b 0D tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
shape 1D tensor of index or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.rng_bit_generator (mhlo::RngBitGeneratorOp)

RngBitGenerator operation

Returns an output filled with uniform random data and an updated output state output_state given an initial state initial_state using the pseudorandom number generator algorithm rng_algorithm .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rng_bit_generator

Пример:

%output_state, %output = mhlo.rng_bit_generator %initial_state, algorithm = THREE_FRY : (tensor<2xui64>) -> (tensor<2xui64>, tensor<2x2xui64>)

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
rng_algorithm ::mlir::mhlo::RngAlgorithmAttr XLA PRNG algorithm to be used.

Operands:

Операнд Описание
initial_state ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Results:

Результат Описание
output_state ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
output statically shaped tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.round_nearest_afz (mhlo::RoundOp)

Round operation

Синтаксис:

operation ::= `mhlo.round_nearest_afz` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise rounding towards the nearest integer, breaking ties away from zero, on the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#round_nearest_afz

Пример:

%result = mhlo.round_nearest_afz %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.round_nearest_even (mhlo::RoundNearestEvenOp)

RoundNearestEven operation

Синтаксис:

operation ::= `mhlo.round_nearest_even` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise rounding towards the nearest integer, breaking ties towards the even integer, on the operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#round_nearest_even

Пример:

%result = mhlo.round_nearest_even %operand : tensor<5xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.rsqrt (mhlo::RsqrtOp)

Rsqrt operation

Синтаксис:

operation ::= `mhlo.rsqrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise reciprocal square root operation on operand tensor and produces a result tensor, implementing the rSqrt operation from the IEEE-754 specification.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#rsqrt

Пример:

%result = mhlo.rsqrt %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.scatter (mhlo::ScatterOp)

Scatter operation

Produces results tensors which are equal to inputs tensors except that several slices specified by scatter_indices are updated with the values updates using update_computation .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#scatter

Пример:

%result = "mhlo.scatter"(%input, %scatter_indices, %update) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = mhlo.add %arg0, %arg1 : tensor<i32>
    mhlo.return %0 : tensor<i32>
}) {
  scatter_dimension_numbers = #mhlo.scatter<
    update_window_dims = [3, 4],
    inserted_window_dims = [1],
    input_batching_dims = [0],
    scatter_indices_batching_dims = [1],
    scatter_dims_to_operand_dims = [2, 1],
    index_vector_dim = 3>,
  indices_are_sorted = false,
  unique_indices = false
} : (tensor<2x3x4x2xi64>, tensor<2x2x3x2xi64>, tensor<2x2x3x2x2xi64>) -> tensor<2x3x4x2xi64>

Traits: RecursiveMemoryEffects , SameVariadicOperandSize

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
scatter_dimension_numbers ::mlir::mhlo::ScatterDimensionNumbersAttr Attribute that models the dimension information for scatter
indices_are_sorted ::mlir::BoolAttr bool attribute
unique_indices ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
scatter_indices ranked tensor of integer or index values
updates variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.select (mhlo::SelectOp)

Select operation

Синтаксис:

operation ::= `mhlo.select` operands attr-dict `:`
              custom<SelectOpType>(type($pred), type($on_true), type($on_false), type($result))

Produces a result tensor where each element is selected from on_true or on_false tensor based on the value of the corresponding element of pred .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#select

Пример:

%result = mhlo.select %pred, %on_true, %on_false : tensor<2x2xi1>, tensor<2x2xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_BroadcastingElementwise , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
pred ranked tensor of pred (AKA boolean or 1-bit integer) values
on_true ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
on_false ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.select_and_scatter (mhlo::SelectAndScatterOp)

SelectAndScatter operation

Scatters the values from the source tensor using scatter based on the outcome of reduce_window of the input tensor using select and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#select_and_scatter

Пример:

%result = "mhlo.select_and_scatter"(%operand, %source, %init_value) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GE>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %arg1) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.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>

Traits: RecursiveMemoryEffects

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
window_dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
window_strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
padding ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
source ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
init_value ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.send (mhlo::SendOp)

Send operation

Sends inputs to a channel channel_id and produces a result token.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#send

Пример:

%result = "mhlo.send"(%operand, %token) {
  // channel_id = 5 : i64,
  // channel_type = #stablehlo<channel_type DEVICE_TO_HOST>,
  channel_handle = #mhlo.channel_handle<handle = 5, type = 2>,
  is_host_transfer = true
} : (tensor<3x4xi32>, !mhlo.token) -> !mhlo.token

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
token жетон

Results:

Результат Описание
«unnamed» жетон

mhlo.set_dimension_size (mhlo::SetDimensionSizeOp)

SetDimensionSize operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/8

Informally, this operation does the same thing as XLA's SetDimensionSize: https://www.tensorflow.org/xla/operation_semantics#setdimensionsize

Пример:

%0 = mhlo.set_dimension_size %arg0, %arg1, dim = 1 : (tensor<4x2xf32>, tensor<i32>) -> tensor<4x2xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
size tensor of 32-bit signless integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.shift_left (mhlo::ShiftLeftOp)

ShiftLeft operation

Синтаксис:

operation ::= `mhlo.shift_left` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise left-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_left

Пример:

%result = mhlo.shift_left %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.shift_right_arithmetic (mhlo::ShiftRightArithmeticOp)

ShiftRightArithmetic operation

Синтаксис:

operation ::= `mhlo.shift_right_arithmetic` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise arithmetic right-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_right_arithmetic

Пример:

%result = mhlo.shift_right_arithmetic %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.shift_right_logical (mhlo::ShiftRightLogicalOp)

ShiftRightLogical operation

Синтаксис:

operation ::= `mhlo.shift_right_logical` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise logical right-shift operation on the lhs tensor by rhs number of bits and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#shift_right_logical

Пример:

%result = mhlo.shift_right_logical %lhs, %rhs : tensor<6xi8>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

mhlo.sign (mhlo::SignOp)

Sign operation

Синтаксис:

operation ::= `mhlo.sign` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Returns the sign of the operand element-wise and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sign

Пример:

%result = mhlo.sign %operand : tensor<7xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.sine (mhlo::SineOp)

Sine operation

Синтаксис:

operation ::= `mhlo.sine` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise sine operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sine

Пример:

%result = mhlo.sine %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.slice (mhlo::SliceOp)

Slice operation

Extracts a slice from the operand using statically-computed starting indices and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#slice

Пример:

%result = "mhlo.slice" (%operand) {
  start_indices = dense<[1, 2]> : tensor<2xi64>,
  limit_indices = dense<[3, 4]> : tensor<2xi64>,
  strides = dense<1> : tensor<2xi64>
} : (tensor<3x4xi64>) -> tensor<2x2xi64>

Traits: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
start_indices ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
limit_indices ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute
strides ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sort (mhlo::SortOp)

Sort operation

Sorts a variadic number of tensors in inputs together, according to a custom comparator , along the given dimension and produces a variadic number of tensors as results .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sort

Пример:

%result0, %result1 = "mhlo.sort"(%input0, %input1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>, %arg2: tensor<i32>, %arg3: tensor<i32>):
    %predicate = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction GT>
      } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%predicate) : (tensor<i1>) -> ()
}) {
  dimension = 0 : i64,
  is_stable = true
} : (tensor<2x3xi32>, tensor<2x3xi32>) -> (tensor<2x3xi32>, tensor<2x3xi32>)

Traits: InferTensorType , RecursiveMemoryEffects , SameOperandsAndResultShape

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
dimension ::mlir::IntegerAttr 64-bit signless integer attribute
is_stable ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
inputs variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sparse_dot (mhlo::SparseDotOp)

Sparse dot operation

Similar to dot_general operation, with one or both of the operands being sparse. An additional argument provides sparsity meta information. Disclaimer: this op is experimental / a work in progress.

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
lhs_sparsity ::mlir::mhlo::SparsityDescriptorAttr Describes structured (N:M) sparsity configuration
rhs_sparsity ::mlir::mhlo::SparsityDescriptorAttr Describes structured (N:M) sparsity configuration
dot_dimension_numbers ::mlir::mhlo::DotDimensionNumbersAttr Attribute that models the dimension information for dot.
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Операнд Описание
lhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
meta variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.sqrt (mhlo::SqrtOp)

Sqrt operation

Синтаксис:

operation ::= `mhlo.sqrt` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise square root operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#sqrt

Пример:

%result = mhlo.sqrt %operand : tensor<2x2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.stochastic_convert (mhlo::StochasticConvertOp)

StochasticConvert operation

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/295

Informally, this operation performs element-wise conversion of values from a bigger type to a smaller one with stochastic rounding using the random number passed in.

Traits: AlwaysSpeculatableImplTrait , Elementwise

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values
random ranked tensor of 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.subtract (mhlo::SubtractOp)

Subtract operation

Синтаксис:

operation ::= `mhlo.subtract` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise subtraction of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#subtract

Пример:

%result = mhlo.subtract %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
rhs ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.tan (mhlo::TanOp)

Tan operation

Синтаксис:

operation ::= `mhlo.tan` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

This operation is a work in progress, so it is not yet included in the specification: https://github.com/openxla/stablehlo/issues/954

Informally, this operation returns Tan(operand) element-wise.

Пример:

%0 = mhlo.tan %arg0 : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.tanh (mhlo::TanhOp)

Tanh operation

Синтаксис:

operation ::= `mhlo.tanh` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise hyperbolic tangent operation on operand tensor and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#tanh

Пример:

%result = mhlo.tanh %operand : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values

mhlo.topk (mhlo::TopKOp)

TopK operation

Синтаксис:

operation ::= `mhlo.topk` `(`$operand `,` `k` `=` $k (`,` `largest` `=` $largest^)? `)` attr-dict `:`
              type($operand) `->` `(`type($values)`,` type($indices)`)`

Returns top k values and their indices, along the last dimension of the operand if largest=true or the bottom k values if largest=false .

See: https://www.tensorflow.org/xla/operation_semantics#top-k

Пример:

%values, %indices = mhlo.topk(%operand, k=5, largest=true)
  : tensor<100xf32> -> (tensor<5xf32>, tensor<5xi32>)

Traits: InferTensorType , RecursiveMemoryEffects

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
k ::mlir::IntegerAttr 64-bit signless integer attribute
largest ::mlir::BoolAttr bool attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
values ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
indices ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.torch_index_select (mhlo::TorchIndexSelectOp)

TorchIndexSelect operation

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/3

Informally, this operation does the same thing as PyTorch's index_select, augmented with support for batch dimensions: https://pytorch.org/docs/stable/generated/torch.index_select.html

The batch_dims attribute specifies the number of major batch dimensions (0 or more) that act like a multidimensional loop over both the operand and the index.

Пример:

%result = "mhlo.torch_index_select"(%operand, %index) {
  dim = 2 : i64,
  batch_dims = 1 : i64
} : (tensor<8x128x3072x64xf32>, tensor<8x16x1024xi32>) -> tensor<8x128x16x1024x64xf32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
dim ::mlir::IntegerAttr 64-bit signless integer attribute
batch_dims ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values
index ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.trace (mhlo::TraceOp)

Trace operation

Синтаксис:

operation ::= `mhlo.trace` $operand `,` $tag attr-dict `:` type($operand)

This operation is on its way out of StableHLO, so it is not included in the specification: https://github.com/openxla/stablehlo/issues/604

It is not used by JAX, PyTorch or TensorFlow, so it looks like we should've classified it as "Private to XLA" and not included it in StableHLO in the first place. With that in mind, its semantics will not be documented here.

Пример:

mhlo.trace %arg0, "In test code." : tensor<5x1x5xi32>

Attributes:

Атрибут MLIR Type Описание
tag ::mlir::StringAttr string attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.transpose (mhlo::TransposeOp)

Transpose operation

Permutes the dimensions of operand tensor using permutation and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#transpose

Пример:

%0 = mhlo.transpose %arg0, dims = [2, 1, 0] : (tensor<1x2x3xi32>) -> tensor<3x2x1xi32>

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
permutation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.triangular_solve (mhlo::TriangularSolveOp)

TriangularSolve operation

Solves batches of systems of linear equations with lower or upper triangular coefficient matrices.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#triangular_solve

Пример:

%result = "mhlo.triangular_solve"(%a, %b) {
  left_side = true,
  lower = true,
  unit_diagonal = false,
  transpose_a = #stablehlo<transpose NO_TRANSPOSE>
} : (tensor<3x3xf32>, tensor<3x3xf32>) -> tensor<3x3xf32>

Traits: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Атрибут MLIR Type Описание
left_side ::mlir::BoolAttr bool attribute
lower ::mlir::BoolAttr bool attribute
unit_diagonal ::mlir::BoolAttr bool attribute
transpose_a ::mlir::mhlo::TransposeAttr Transpose options

Operands:

Операнд Описание
a ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values
b ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

Results:

Результат Описание
«unnamed» ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or complex type with 32-bit float or 64-bit float elements values

mhlo.tuple (mhlo::TupleOp)

Tuple operation

Синтаксис:

operation ::= `mhlo.tuple` $val attr-dict `:` custom<TupleOpType>(type($val), type($result))

Produces a result tuple from values val .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#tuple

Пример:

%result = mhlo.tuple %val0, %val1 : tuple<tensor<2xf32>, tuple<tensor<i32>>>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
val variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token or nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

Results:

Результат Описание
result nested tuple with any combination of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token values

mhlo.uniform_dequantize (mhlo::UniformDequantizeOp)

UniformDequantize operation

Синтаксис:

operation ::= `mhlo.uniform_dequantize` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise conversion of quantized tensor operand to a floating-point tensor result according to the quantization parameters defined by the operand type.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#uniform_dequantize

Пример:

%result = mhlo.uniform_dequantize %operand : (tensor<16x16x!quant.uniform<i8:f32, 34.0:16>>) -> tensor<16x16xf32>

Traits: AlwaysSpeculatableImplTrait , Elementwise , InferTensorType , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type values

mhlo.uniform_quantize (mhlo::UniformQuantizeOp)

UniformQuantize operation

Синтаксис:

operation ::= `mhlo.uniform_quantize` $operand attr-dict
              `:` custom<SameOperandsAndResultType>(type($operand), type($result))

Performs element-wise conversion of floating-point tensor or quantized tensor operand to a quantized tensor result according to the quantization parameters defined by the result type.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#uniform_quantize

Пример:

%result = mhlo.uniform_quantize %operand : (tensor<16x16xf32>) -> tensor<16x16x!quant.uniform<ui8:f32, 34.0:16>>

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
operand ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Результаты:

Результат Описание
result ranked tensor of 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

mhlo.while (mhlo::WhileOp)

While operation

Produces the output from executing body function 0 or more times while the cond function outputs true .

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#while

Пример:

%results0, %results1 = "mhlo.while"(%operand0, %operand1) ({
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.compare"(%arg0, %arg1) {
      comparison_direction = #stablehlo<comparison_direction LT>
    } : (tensor<i32>, tensor<i32>) -> tensor<i1>
    "mhlo.return"(%0) : (tensor<i1>) -> ()
}, {
  ^bb0(%arg0: tensor<i32>, %arg1: tensor<i32>):
    %0 = "mhlo.add"(%arg0, %constant0) : (tensor<i32>, tensor<i32>) -> tensor<i32>
    "mhlo.return"(%0, %arg1) : (tensor<i32>, tensor<i32>) -> ()
}) : (tensor<i32>, tensor<i32>) -> (tensor<i32>, tensor<i32>)

Traits: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface , OpAsmOpInterface

Operands:

Операнд Описание
operand variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

Results:

Результат Описание
«unnamed» variadic of ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer values or ranked tensor of 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values or token

mhlo.xla.rng_get_and_update_state (mhlo::XlaRngGetAndUpdateStateOp)

XlaRngGetAndUpdateState operation

Синтаксис:

operation ::= `mhlo.xla.rng_get_and_update_state` attr-dict

This operation is private to the XLA compiler, so it is does not yet have a specification.

Informally, this operation represents the change of the global random number generator state for rng instructions. The global state is incremented by delta and the old state is returned.

The output is currently defined for a single output type. If this changes in the future to support multiple types, lowering to use of a global memref must ensure that a single memref is still used and updated appropriately.

Interfaces: InferTypeOpInterface

Attributes:

Атрибут MLIR Type Описание
delta ::mlir::IntegerAttr 64-bit signless integer attribute

Results:

Результат Описание
«unnamed» statically shaped tensor of 64-bit unsigned integer values

mhlo.xor (mhlo::XorOp)

Xor operation

Синтаксис:

operation ::= `mhlo.xor` $lhs `,` $rhs attr-dict
              `:` custom<SameOperandsAndResultType>(type($lhs), type($rhs), type($result))

Performs element-wise XOR of two tensors lhs and rhs and produces a result tensor.

See: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#xor

Пример:

%result = mhlo.xor %lhs, %rhs : tensor<2xi32>

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Операнд Описание
lhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values
rhs ranked tensor of pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer values

Results:

Результат Описание
result ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/8/16/32-bit uniform quantized per axis signed integer or 2/4/8/16/32-bit uniform quantized per axis unsigned integer values

Атрибуты

ArgResultAliasAttr

Attribute that models the alias relationship of entry function argument

This attribute captures the alias relationship of an MHLO main function argument to one of the results, denoted by resultIndex . The argTupleIndices and resultTupleIndices are used to index into nested tuples in operand and result respectively. If isMustAlias is true then the operand-result pair must alias.

This is meant to be used as an attribute on a function argument in MHLO. For example, in the following code it expresses that %arg1 may alias 0-th result.

func @main(%arg0: tensor<2xf32>, %arg1: tensor<3xf32> {mhlo.result_alias =
    mhlo.result_alias<result_index = [2], ...>}
  ) -> tensor<2xf32>, tensor<3xf32> {
  // function body ...
}

Параметры:

Параметр C++ type Описание
argTupleIndices ::llvm::ArrayRef<int64_t> Измерение
resultIndex int64_t
resultTupleIndices ::llvm::ArrayRef<int64_t> Измерение
isMustAlias bool

ChannelHandleAttr

two 64-bit integers 'handle' and 'type'

Синтаксис:

#mhlo.channel_handle<
  int64_t,   # handle
  int64_t   # type
>

Параметры:

Параметр C++ type Описание
ручка int64_t
тип int64_t

ComparisonDirectionAttr

Which comparison operation to perform.

Синтаксис:

#mhlo.comparison_direction<
  ::mlir::mhlo::ComparisonDirection   # value
>

Enum cases:

  • EQ ( EQ )
  • NE ( NE )
  • GE ( GE )
  • GT ( GT )
  • LE ( LE )
  • LT ( LT )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::ComparisonDirection an enum of type ComparisonDirection

ComparisonTypeAttr

Which comparison type to use.

Синтаксис:

#mhlo.comparison_type<
  ::mlir::mhlo::ComparisonType   # value
>

Enum cases:

  • NOTYPE ( NOTYPE )
  • FLOAT ( FLOAT )
  • TOTALORDER ( TOTALORDER )
  • SIGNED ( SIGNED )
  • UNSIGNED ( UNSIGNED )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::ComparisonType an enum of type ComparisonType

ConvDimensionNumbersAttr

Structure of dimension information for conv op

Параметры:

Параметр C++ type Описание
inputBatchDimension int64_t
inputFeatureDimension int64_t
inputSpatialDimensions ::llvm::ArrayRef<int64_t> Измерение
kernelInputFeatureDimension int64_t
kernelOutputFeatureDimension int64_t
kernelSpatialDimensions ::llvm::ArrayRef<int64_t> Измерение
outputBatchDimension int64_t
outputFeatureDimension int64_t
outputSpatialDimensions ::llvm::ArrayRef<int64_t> Измерение

CrossProgramPrefetchAttr

Argument that is prefetched from another program

Синтаксис:

#mhlo.cross_program_prefetch<
  int64_t,   # parameter
  ::llvm::ArrayRef<int64_t>,   # indices
  std::optional<int64_t>   # offset
>

This attribute captures an argument that is prefetched from another program. For a given CrossProgramPrefetchAttr , parameter tells us which argument of the main function of the module is prefetched, and indices is a shape index telling us what subshape of that argument is prefetched.

A shape has a subshape iff it is a tuple. In that case, the subshape of the tuple by indices is the shape achieved after indexing by each element of indices in turn. For example, the [1,0] subshape of tuple<tuple<token, token>, tuple<tensor<i32>, token>> is tensor<i32> .

An empty value for indices means the whole shape is prefetched.

Например,

module attributes { mhlo.cross_program_prefetch = [ #mhlo.cross_program_prefetch< parameter = 0, indices = [0]> ]} {
  func.func @copy(%arg0 : tuple<tensor<2x3xi32>, tensor<i32>>) -> tuple<tensor<2x3xi32>, tensor<i32>> {
    %0 = "mhlo.copy"(%arg0) {is_cross_program_prefetch}
    return %0 : tuple<tensor<2x3xi32>, tensor<i32>>
  }
  func.func @main(%arg0 : tuple<tensor<2x3xi32>, tensor<i32>>) -> tuple<tensor<2x3xi32>, tensor<i32>> {
    %1 = "mhlo.async_start"(%arg0) {called_computation=@copy}
    %2 = "mhlo.async_done"(%1) {called_computation=@copy}
    return %2 : tuple<tensor<2x3xi32>, tensor<i32>>
  }
}

The parameter = 0 tells us that the async copy of the 0 th parameter is a cross_program_prefetch , while the index of [0] tells us that the 0 th element of the tuple is prefetched while the other element of the tuple is not.

Параметры:

Параметр C++ type Описание
параметр int64_t
индексы ::llvm::ArrayRef<int64_t> Измерение
компенсировать std::optional<int64_t>

CustomCallScheduleAttr

Specifies the desired schedule for the custom-call.

Синтаксис:

#mhlo.custom_call_schedule<
  ::mlir::mhlo::CustomCallSchedule   # value
>

Enum cases:

  • NONE ( NONE )
  • LATEST ( LATEST )
  • EARLIEST ( EARLIEST )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::CustomCallSchedule an enum of type CustomCallSchedule

DequantizeModeAttr

Dequantization mode. Only MIN_COMBINED is supported.

Синтаксис:

#mhlo.dequantize_mode<
  ::mlir::mhlo::DequantizeMode   # value
>

Enum cases:

  • MIN_COMBINED ( MIN_COMBINED )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::DequantizeMode an enum of type DequantizeMode

DomainKindAttr

Kind of domain metatdata attached to an HLO domain.

Синтаксис:

#mhlo.kind<
  ::mlir::mhlo::DomainKind   # value
>

Enum cases:

  • sharding ( sharding )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::DomainKind an enum of type DomainKind

DotAlgorithmAttr

Attribute that models the algorithm constraints to use for computing dot.

Синтаксис:

#mhlo.dot_algorithm<
  Type,   # lhsPrecisionType
  Type,   # rhsPrecisionType
  Type,   # accumulationType
  int64_t,   # lhsComponentCount
  int64_t,   # rhsComponentCount
  int64_t,   # numPrimitiveOperations
  bool   # allowImpreciseAccumulation
>

Параметры:

Параметр C++ type Описание
lhsPrecisionType Type
rhsPrecisionType Type
accumulationType Type
lhsComponentCount int64_t
rhsComponentCount int64_t
numPrimitiveOperations int64_t
allowImpreciseAccumulation bool

DotDimensionNumbersAttr

Attribute that models the dimension information for dot.

Параметры:

Параметр C++ type Описание
lhsBatchingDimensions ::llvm::ArrayRef<int64_t> Измерение
rhsBatchingDimensions ::llvm::ArrayRef<int64_t> Измерение
lhsContractingDimensions ::llvm::ArrayRef<int64_t> Измерение
rhsContractingDimensions ::llvm::ArrayRef<int64_t> Измерение

FftTypeAttr

XLA fast fourier transform type.

Синтаксис:

#mhlo.fft_type<
  ::mlir::mhlo::FftType   # value
>

Enum cases:

  • FFT ( FFT )
  • IFFT ( IFFT )
  • RFFT ( RFFT )
  • IRFFT ( IRFFT )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::FftType an enum of type FftType

FusionKindAttr

fusion kind

Синтаксис:

#mhlo.fusion_kind<
  ::mlir::mhlo::FusionKind   # value
>

Enum cases:

  • kLoop ( kLoop )
  • kInput ( kInput )
  • kOutput ( kOutput )
  • kCustom ( kCustom )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::FusionKind an enum of type FusionKind

GatherDimensionNumbersAttr

Attribute that models the dimension information for gather

Параметры:

Параметр C++ type Описание
offsetDims ::llvm::ArrayRef<int64_t> Измерение
collapsedSliceDims ::llvm::ArrayRef<int64_t> Измерение
operandBatchingDims ::llvm::ArrayRef<int64_t> Измерение
startIndicesBatchingDims ::llvm::ArrayRef<int64_t> Измерение
startIndexMap ::llvm::ArrayRef<int64_t> Измерение
indexVectorDim int64_t

OutputOperandAliasAttr

Attribute that models the alias relationship of output and operand of a CustomCall op

Синтаксис:

#mhlo.output_operand_alias<
  ::llvm::ArrayRef<int64_t>,   # outputTupleIndices
  int64_t,   # operandIndex
  ::llvm::ArrayRef<int64_t>   # operandTupleIndices
>

This attribute captures the alias relationship of the output to one of the operands for a CustomCall op, denoted by operand_index . The output_tuple_indices and operand_tuple_indices are used to index into output and operand types. These indices lists are empty if the corresponding types are not tuple types, and can be arbitrarily long in case of arbitrarily nested tuple types.

See https://www.tensorflow.org/xla/aliasing

Example when used as array with in mhlo.custom-call:

%0 = "mhlo.custom_call"(%arg0, %arg1) {
  // other attributes
  output_operand_alias = [
    #mhlo.output_operand_alias<output_tuple_indices = [0],
                               operand_index = 0,
                               operand_tuple_indices = [1]>
  ]
} : (tuple<tensor<1x1xf32>, tensor<2x3xf32>>, tensor<5x5xf32>) -> tuple<tensor<2x3xf32>>

The output and the 0th operand are both tuples. The aliasing shows the
relationship between the 0th element in output tuple with the 1st element in
the 0th operand. And both of them are of the same type: tensor<2x3xf32>.

Параметры:

Параметр C++ type Описание
outputTupleIndices ::llvm::ArrayRef<int64_t> Измерение
operandIndex int64_t
operandTupleIndices ::llvm::ArrayRef<int64_t> Измерение

PrecisionAttr

XLA precision for an operand. Has backend specific meaning.

Синтаксис:

#mhlo.precision<
  ::mlir::mhlo::Precision   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • HIGH ( HIGH )
  • HIGHEST ( HIGHEST )
  • PACKED_NIBBLE ( PACKED_NIBBLE )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::Precision an enum of type Precision

RngAlgorithmAttr

XLA PRNG algorithm to be used.

Синтаксис:

#mhlo.rng_algorithm<
  ::mlir::mhlo::RngAlgorithm   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • THREE_FRY ( THREE_FRY )
  • PHILOX ( PHILOX )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::RngAlgorithm an enum of type RngAlgorithm

RngDistributionAttr

XLA PRNG distribution to be used.

Синтаксис:

#mhlo.rng_distribution<
  ::mlir::mhlo::RngDistribution   # value
>

Enum cases:

  • UNIFORM ( UNIFORM )
  • NORMAL ( NORMAL )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::RngDistribution an enum of type RngDistribution

ScatterDimensionNumbersAttr

Attribute that models the dimension information for scatter

Параметры:

Параметр C++ type Описание
updateWindowDims ::llvm::ArrayRef<int64_t> Измерение
insertedWindowDims ::llvm::ArrayRef<int64_t> Измерение
inputBatchingDims ::llvm::ArrayRef<int64_t> Измерение
scatterIndicesBatchingDims ::llvm::ArrayRef<int64_t> Измерение
scatterDimsToOperandDims ::llvm::ArrayRef<int64_t> Измерение
indexVectorDim int64_t

SparsityDescriptorAttr

Describes structured (N:M) sparsity configuration

Синтаксис:

#mhlo.sparsity<
  int64_t,   # dimension
  int64_t,   # n
  int64_t   # m
>

This attribute is defined for a sparse dot operation with a structured sparse input tensor. With (N=2,M=4), every 4 consecutive logical elements have exactly 2 non-zero physical elements in the input tensor.

$dimension defines the index of the contracting dimension that is sparse (it has to be the most minor dimension). The additional metadata operand in the sparse dot operation defines which logical elements are zeroed out.

Параметры:

Параметр C++ type Описание
измерение int64_t
н int64_t
м int64_t

TransposeAttr

Transpose options

Синтаксис:

#mhlo.transpose<
  ::mlir::mhlo::Transpose   # value
>

Enum cases:

  • TRANSPOSE_INVALID ( TRANSPOSE_INVALID )
  • NO_TRANSPOSE ( NO_TRANSPOSE )
  • TRANSPOSE ( TRANSPOSE )
  • ADJOINT ( ADJOINT )

Параметры:

Параметр C++ type Описание
ценить ::mlir::mhlo::Transpose an enum of type Transpose

TypeExtensionsAttr

Attribute that extends tensor type with MHLO type properties.

Синтаксис:

#mhlo.type_extensions<
  ::llvm::ArrayRef<int64_t>   # bounds
>

This attribute is used to extend MLIR tensor type with MHLO tensor specific properties. These properties aren't modeled in the MLIR type. This attribute is set in the encoding field of the tensor type.

See HLO_BoundedAttrInterface for documentation for bounds .

Параметры:

Параметр C++ type Описание
границы ::llvm::ArrayRef<int64_t>

Типы

AsyncBundleType

Opaque collection of other types

Синтаксис:

!mhlo.async_bundle<
  ::llvm::ArrayRef<Type>   # types
>

Параметры:

Параметр C++ type Описание
типы ::llvm::ArrayRef<Type>

Перечисления

ComparisonDirection

Which comparison operation to perform.

Случаи:

Символ Ценить Нить
эквалайзер 0 эквалайзер
СВ 1 СВ
GE 2 GE
ГТ 3 ГТ
ЛЕ 4 ЛЕ
LT 5 LT

ComparisonType

Which comparison type to use.

Случаи:

Символ Ценить Нить
NOTYPE 0 NOTYPE
ПЛАВАТЬ 1 ПЛАВАТЬ
TOTALORDER 2 TOTALORDER
ПОДПИСАНО 3 ПОДПИСАНО
UNSIGNED 4 UNSIGNED

CustomCallApiVersion

Custom call API version

Случаи:

Символ Ценить Нить
API_VERSION_UNSPECIFIED 0 API_VERSION_UNSPECIFIED
API_VERSION_ORIGINAL 1 API_VERSION_ORIGINAL
API_VERSION_STATUS_RETURNING 2 API_VERSION_STATUS_RETURNING
API_VERSION_STATUS_RETURNING_UNIFIED 3 API_VERSION_STATUS_RETURNING_UNIFIED
API_VERSION_TYPED_FFI 4 API_VERSION_TYPED_FFI

CustomCallSchedule

Specifies the desired schedule for the custom-call.

Случаи:

Символ Ценить Нить
НИКТО 0 НИКТО
ПОСЛЕДНИЙ 1 ПОСЛЕДНИЙ
EARLIEST 2 EARLIEST

DequantizeMode

Dequantization mode. Only MIN_COMBINED is supported.

Случаи:

Символ Ценить Нить
MIN_COMBINED 0 MIN_COMBINED

DomainKind

Kind of domain metatdata attached to an HLO domain.

Случаи:

Символ Ценить Нить
шардинг 0 шардинг

FftType

XLA fast fourier transform type.

Случаи:

Символ Ценить Нить
БПФ 0 БПФ
IFFT 1 IFFT
RFFT 2 RFFT
IRFFT 3 IRFFT

FusionKind

fusion kind

Случаи:

Символ Ценить Нить
kLoop 0 kLoop
kInput 1 kInput
kOutput 2 kOutput
kCustom 3 kCustom

Точность

XLA precision for an operand. Has backend specific meaning.

Случаи:

Символ Ценить Нить
ПО УМОЛЧАНИЮ 0 ПО УМОЛЧАНИЮ
ВЫСОКИЙ 1 ВЫСОКИЙ
HIGHEST 2 HIGHEST
PACKED_NIBBLE 3 PACKED_NIBBLE

RngAlgorithm

XLA PRNG algorithm to be used.

Случаи:

Символ Ценить Нить
ПО УМОЛЧАНИЮ 0 ПО УМОЛЧАНИЮ
THREE_FRY 1 THREE_FRY
PHILOX 2 PHILOX

RngDistribution

XLA PRNG distribution to be used.

Случаи:

Символ Ценить Нить
UNIFORM 1 UNIFORM
НОРМАЛЬНЫЙ 2 НОРМАЛЬНЫЙ

Transpose

Transpose options

Случаи:

Символ Ценить Нить
TRANSPOSE_INVALID 0 TRANSPOSE_INVALID
NO_TRANSPOSE 1 NO_TRANSPOSE
TRANSPOSE 2 TRANSPOSE
ADJOINT 3 ADJOINT