'mhlo' Dialeto

Operações

mhlo.abs (mhlo::AbsOp)

Operação abdominal

Sintaxe:

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

Executa operação abs elemento a elemento no tensor operand e produz um tensor result .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Efeitos: MemoryEffects::Effect{}

Operandos:

Operando Descrição
operand tensor classificado de número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou 0FNU tipo ou float de 16 bits ou float de 32 bits ou float de 64 bits ou tipo bfloat16 ou tipo complexo com float de 32 bits ou elementos float de 64 bits ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado inteiro sem sinal ou 2/4/8/16/32 bits uniforme quantizado por eixo sem sinal valores inteiros

Resultados:

Resultado Descrição
result Tensor classificado de 2/4/8/16/32/64 bits sem sinais ou tipo f4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou tipo f8e4 ou f8e3M3 Nuz tipo ou F8E8M0FNU tipo ou float de 16 bits ou float de 32 bits ou float de 64 bits ou tipo bfloat16 ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou quantizado uniforme de 2/4/8/16/32 bits por inteiro com sinal do eixo ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou valores inteiros sem sinal quantizados uniformes de 2/4/8/16/32 bits por eixo

mhlo.add (mhlo :: addop)

Adicionar operação

Sintaxe:

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

Executa a adição elemento a elemento de dois tensores lhs e rhs e produz um tensor result .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Efeitos: MemoryEffects::Effect{}

Operandos:

Operando Descrição
lhs tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com Elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8/16 / Quantizado uniforme de 32 bits por inteiro com sinal de eixo ou 2/4/8/16/32 bits quantizado uniforme por valores inteiros sem sinal de eixo
rhs tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com Elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8/16 / Quantizado uniforme de 32 bits por inteiro com sinal de eixo ou 2/4/8/16/32 bits quantizado uniforme por valores inteiros sem sinal de eixo

Resultados:

Resultado Descrição
result tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com Elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8/16 / Quantizado uniforme de 32 bits por inteiro com sinal de eixo ou 2/4/8/16/32 bits quantizado uniforme por valores inteiros sem sinal de eixo

mhlo.add_dependency (mhlo::AddDependencyOp)

Operação AdicionarDependência

Sintaxe:

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

Esta operação é privada do compilador XLA, portanto ainda não possui especificação.

Informalmente, esta operação tem dois operandos: um operando de dados e um token. A saída da operação é o operando de dados. Quando usada com AfterAll, esta operação permite ordenar operações sem efeitos colaterais (aquelas que não produzem valores de token).

Exemplo:

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

Características: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Efeitos: MemoryEffects::Effect{}

Operandos:

Operando Descrição
operand tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com Elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiros assinados quantizados uniformes de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits ou tensor classificado de 2/4 /8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou token
token ficha

Resultados:

Resultado Descrição
output tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com Elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiros assinados quantizados uniformes de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits ou tensor classificado de 2/4 /8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou token

mhlo.after_all (mhlo::AfterAllOp)

Operação AfterAll

Sintaxe:

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

Garante que as operações que produzem as inputs sejam executadas antes de quaisquer operações que dependam de result .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Efeitos: MemoryEffects::Effect{}

Operandos:

Operando Descrição
inputs variável de token

Resultados:

Resultado Descrição
result ficha

mhlo.all_gather (mhlo::AllGatherOp)

Operação AllGather

Dentro de cada grupo de processos na grade de processos, concatena os valores do tensor de operando de cada processo ao longo de all_gather_dim e produz um tensor de resultado. O computation é aplicado separadamente para cada operando em operands , produzindo um resultado por operando.

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

Exemplo:

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

Características: SameOperandsAndResultElementType

Atributos:

Atributo Tipo MLIR Descrição
all_gather_dim ::mlir::IntegerAttr Atributo inteiro sem sinal de 64 bits cujo valor não é negativo
replica_groups ::mlir::DenseIntElementsAttr Atributo de elementos inteiros sem sinal de 64 bits
channel_handle ::mlir::mhlo::ChannelHandleAttr dois inteiros de 64 bits 'handle' e 'type'
use_global_device_ids ::mlir::UnitAttr atributo de unidade

Operandos:

Operando Descrição
operands Variada do tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou tipo f8e4m3fn ou tipo f8e4m3fnu ou tipo f8e4m3b11fnUz ou f8e4m3fzz ou tipo f8e4m3b1fnUz ou f8e4m3 Flutuação de 16 bits ou bóia de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro com sinal de eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal de eixo

Resultados:

Resultado Descrição
«sem nome» variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro com sinal de eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal de eixo

mhlo.all_reduce (mhlo::AllReduceOp)

Operação AllReduce

Dentro de cada grupo de processos na grade de processos, aplica um computation de função de redução aos valores de um tensor de operando de cada processo e produz um tensor de resultado. O computation é aplicado separadamente para cada operando em operands , produzindo um resultado por operando.

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

Exemplo:

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

Características: InferTensorType , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Atributos:

Atributo Tipo MLIR Descrição
replica_groups ::mlir::DenseIntElementsAttr Atributo de elementos inteiros sem sinal de 64 bits
channel_handle ::mlir::mhlo::ChannelHandleAttr dois inteiros de 64 bits 'handle' e 'type'
use_global_device_ids ::mlir::UnitAttr atributo de unidade

Operandos:

Operando Descrição
operands variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro com sinal de eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal de eixo

Resultados:

Resultado Descrição
«sem nome» variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro com sinal de eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal de eixo

mhlo.all_to_all (mhlo::AllToAllOp)

Operação AllToAll

Dentro de cada grupo de processos na grade de processos, divide os valores do tensor operand ao longo de split_dimension em partes, espalha as partes divididas entre os processos, concatena as partes dispersas ao longo concat_dimension e produz um tensor result .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsElementType , SameOperandsShape , SameVariadicOperandSize

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo MLIR Descrição
split_dimension ::mlir::IntegerAttr Atributo inteiro sem sinal de 64 bits cujo valor não é negativo
concat_dimension :: mlir :: integerattr Atributo inteiro sem sinal de 64 bits cujo valor não é negativo
split_count ::mlir::IntegerAttr Atributo inteiro sem sinal de 64 bits cujo valor é positivo
replica_groups ::mlir::DenseIntElementsAttr Atributo de elementos inteiros sem sinal de 64 bits
channel_handle ::mlir::mhlo::ChannelHandleAttr dois inteiros de 64 bits 'handle' e 'type'

Operandos:

Operando Descrição
operand variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit bit ou bfloat16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou número inteiro ou complexo 2/4/8/16/32/64 bits tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro com sinal de eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal de eixo

Resultados:

Resultado Descrição
«sem nome» variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro com sinal de eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal de eixo

mhlo.and (mhlo::AndOp)

E operação

Sintaxe:

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

Executa AND elemento a elemento de dois tensores lhs e rhs e produz um tensor result

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Efeitos: MemoryEffects::Effect{}

Operandos:

Operando Descrição
lhs tensor classificado de pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou valores inteiros sem sinal de 2/4/8/16/32/64 bits
rhs tensor classificado de pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou valores inteiros sem sinal de 2/4/8/16/32/64 bits

Resultados:

Resultado Descrição
result tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou número inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com Elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8/16 / Quantizado uniforme de 32 bits por inteiro com sinal de eixo ou 2/4/8/16/32 bits quantizado uniforme por valores inteiros sem sinal de eixo

mhlo.async_done (mhlo::AsyncDoneOp)

Operação AsyncDone

Esta operação é privada do compilador XLA, portanto ainda não possui especificação.

Informalmente, esta operação é bloqueada até o final de uma computação assíncrona. Ele retorna o resultado final da computação assíncrona.

Consulte a documentação do AsyncStart para obter mais informações.

Interfaces: InferTypeOpInterface

Operandos:

Operando Descrição
bundle async_bundle com qualquer combinação de tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou Flutuador de 32 bits ou float de 64 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou 2/4/8/16/32/64 bits sem sinal tipo inteiro ou complexo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/ 4/8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou valores de token

Resultados:

Resultado Descrição
«sem nome» variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro assinado por eixo ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou token ou tupla aninhada com qualquer combinação de tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou f6E3M2FN tipo ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou float de 64 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou 1 -bit inteiro) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits ou tensor classificado de quantizado uniforme uniforme de 2/4/8/16/32 bits inteiro assinado por eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal ou valores de token por eixo

mhlo.async_start (mhlo::AsyncStartOp)

Operação AsyncStart

Esta operação é privada do compilador XLA, portanto ainda não possui especificação.

Informalmente, esta operação inicia uma computação assíncrona.

Isso é usado quando há funções que contêm esperas assíncronas (como DMAs) e computação no thread. Por exemplo, uma função pode consistir em um cálculo, um DMA, outro cálculo, um segundo DMA e um cálculo final. Isso seria representado como async_start seguido por async_update e async_done. O async_start faria o primeiro cálculo no thread e depois iniciaria o DMA. O async_update esperaria a conclusão do DMA se ainda não tivesse sido feito, então executaria o segundo cálculo na função e iniciaria o segundo DMA. Finalmente, o async_done esperaria neste último DMA e então executaria o último cálculo que precisa ser executado no thread e retornaria o resultado desse cálculo final.

operands são passados ​​​​diretamente para a computação called_computation é a função que será executada de forma assíncrona execution_thread é o nome do thread no qual ela será executada. O thread principal é chamado de "principal". Todos os tópicos têm nomes.

Isso retorna todo o estado necessário entre operações assíncronas. Após a atribuição do buffer, os valores de retorno representam o espaço necessário para armazenar a entrada, os resultados e quaisquer scratchpads necessários ou editados pela operação assíncrona.

Atributos:

Atributo Tipo MLIR Descrição
called_computation ::mlir::FlatSymbolRefAttr atributo de referência de símbolo plano
execution_thread ::mlir::StringAttr atributo de string

Operandos:

Operando Descrição
inputs variável do tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 64- bit float ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou complexo tipo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/4/8 /16/32 bits uniforme quantizado por inteiro assinado por eixo ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou token ou tupla aninhada com qualquer combinação de tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou f6E3M2FN tipo ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou float de 64 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou 1 -bit inteiro) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou inteiro sem sinal de 2/4/8/16/32/64 bits ou tipo complexo com elementos flutuantes de 32 bits ou flutuantes de 64 bits ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits ou tensor classificado de quantizado uniforme uniforme de 2/4/8/16/32 bits inteiro assinado por eixo ou 2/4/8/16/32 bits uniforme quantizado por valores inteiros sem sinal ou valores de token por eixo

Resultados:

Resultado Descrição
«sem nome» async_bundle com qualquer combinação de tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou Flutuador de 32 bits ou float de 64 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou 2/4/8/16/32/64 bits sem sinal tipo inteiro ou complexo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/ 4/8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou valores de token

mhlo.async_update (mhlo::AsyncUpdateOp)

Operação AsyncUpdate

Esta operação é privada do compilador XLA, portanto ainda não possui especificação.

Informalmente, esta operação bloqueia uma computação assíncrona até uma barreira de sincronização. Isso retorna bundle após operar nele.

Consulte a documentação do AsyncStart para obter mais informações.

Interfaces: InferTypeOpInterface

Operandos:

Operando Descrição
bundle async_bundle com qualquer combinação de tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou Flutuador de 32 bits ou float de 64 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou 2/4/8/16/32/64 bits sem sinal tipo inteiro ou complexo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/ 4/8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou valores de token

Resultados:

Resultado Descrição
«sem nome» async_bundle com qualquer combinação de tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou Flutuador de 32 bits ou float de 64 bits ou tipo bfloat16 ou pred (também conhecido como booleano ou inteiro de 1 bit) ou inteiro sem sinal de 2/4/8/16/32/64 bits ou 2/4/8/16/32/64 bits sem sinal tipo inteiro ou complexo com elementos flutuantes de 32 bits ou elementos flutuantes de 64 bits ou inteiro sem sinal quantizado uniforme de 2/4/8/16/32 bits ou inteiro não assinado quantizado uniforme de 2/4/8/16/32 bits ou 2/ 4/8/16/32 bits uniforme quantizado por eixo inteiro assinado ou 2/4/8/16/32 bits uniforme quantizado por eixo valores inteiros sem sinal ou valores de token

mhlo.atan2 (mhlo::Atan2Op)

Operação Atan2

Sintaxe:

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

Executa a operação atan2 elemento a elemento no tensor lhs e rhs e produz um tensor result .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Efeitos: MemoryEffects::Effect{}

Operandos:

Operando Descrição
lhs tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou tipo complexo com elementos float de 32 bits ou float de 64 bits ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits
rhs tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou tipo complexo com elementos float de 32 bits ou float de 64 bits ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits

Resultados:

Resultado Descrição
result tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou tipo bfloat16 ou tipo complexo com elementos float de 32 bits ou float de 64 bits ou inteiro assinado quantizado uniforme de 2/4/8/16/32 bits ou valores inteiros não assinados quantizados uniformes de 2/4/8/16/32 bits

mhlo.batch_norm_grad (mhlo::BatchNormGradOp)

Operação BatchNormGrad

Calcula gradientes de várias entradas de BatchNormTrainingOp retropropagando de grad_output e produz tensores grad_operand , grad_scale e grad_offset .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , InferTensorType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo MLIR Descrição
epsilon ::mlir::FloatAttr Atributo flutuante de 32 bits
feature_index ::mlir::IntegerAttr Atributo inteiro sem sinal de 64 bits cujo valor não é negativo

Operandos:

Operando Descrição
operand tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou valores do tipo bfloat16
scale Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16
mean Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16
variance Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16
grad_output tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou valores do tipo bfloat16

Resultados:

Resultado Descrição
grad_operand tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou valores do tipo bfloat16
grad_scale Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16
grad_offset Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16

mhlo.batch_norm_inference (mhlo::BatchNormInferenceOp)

Operação BatchNormInference

Normaliza o tensor operand em todas as dimensões, exceto na dimensão feature_index , e produz um tensor result .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , InferTensorType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo MLIR Descrição
epsilon ::mlir::FloatAttr Atributo flutuante de 32 bits
feature_index ::mlir::IntegerAttr Atributo inteiro sem sinal de 64 bits cujo valor não é negativo

Operandos:

Operando Descrição
operand tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou valores do tipo bfloat16
scale Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16
offset Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuante de 64 bits ou valores do tipo bfloat16
mean Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16
variance Tensor 1D do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou Flutuador de 64 bits ou valores do tipo bfloat16

Resultados:

Resultado Descrição
result tensor classificado do tipo f4E2M1FN ou tipo f6E2M3FN ou tipo f6E3M2FN ou tipo f8E3M4 ou tipo f8E4M3 ou tipo f8E4M3FN ou tipo f8E4M3FNUZ ou tipo f8E4M3B11FNUZ ou tipo f8E5M2 ou tipo f8E5M2FNUZ ou tipo f8E8M0FNU ou float de 16 bits ou float de 32 bits ou 6 Flutuador de 4 bits ou valores do tipo bfloat16

mhlo.batch_norm_training (mhlo::BatchNormTrainingOp)

Operação BatchNormTraining

Calcula a média e a variação entre dimensões de lote e espaciais e normaliza o tensor operand , para cada recurso na dimensão feature_index e produz tensores output , batch_mean e batch_var .

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

Exemplo:

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

Características: AlwaysSpeculatableImplTrait , InferTensorType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo MLIR Descrição
epsilon ::mlir::FloatAttr Atributo flutuante de 32 bits
feature_index ::mlir::IntegerAttr Atributo inteiro sem sinal de 64 bits cujo valor não é negativo

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou valores do tipo BFLOAT16
scale 1d Tensor do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou valores do tipo BFLOAT16
offset 1d Tensor do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou valores do tipo BFLOAT16

Resultados:

Resultado Descrição
output Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou valores do tipo BFLOAT16
batch_mean 1d Tensor do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou f8er4 ou flutuação de 32 bits ou flutuação de 64 bits ou valores do tipo BFLOAT16
batch_var 1d Tensor do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou valores do tipo BFLOAT16

mhlo.bitcast (mhlo :: bittop)

Operação de bitcast

Sintaxe:

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

Esta operação é privada para o compilador XLA, por isso ainda não possui uma especificação.

Informalmente, esta operação altera a forma da entrada da maneira que o arranjo físico dos elementos permanece inalterado.

Esta operação precisa de informações de layout para entender o "arranjo físico dos elementos", e o suporte ao layout no MHLO é atualmente um trabalho em andamento.

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Efeitos: MemoryEffects::Effect{}

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.bitcast_convert (mhlo :: bitcastconverrop)

Operação BitcastConvert

Sintaxe:

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

Executa uma operação de bitcast no tensor operand e produz um tensor result , onde os bits de todo o tensor operand são reinterpretados usando o tipo do tensor de result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#bitcast_convert

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Efeitos: MemoryEffects::Effect{}

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.broadcast (MHLO :: Broadtop)

Operação de transmissão

Esta operação está saindo de StableHlo, portanto não está incluída na especificação: https://github.com/openxla/stablehlo/issues/3

Informalmente, esta operação faz a mesma coisa que a transmissão do XLA: https://www.tensorflow.org/xla/operation_seMantics#broadcast

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo mlir Descrição
broadcast_sizes :: mlir :: denseInteLentMattr Atributo de elementos inteiros sem sinais de 64 bits

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.broadcast_in_dim (MHLO :: BroadcastIndimop)

Operação Broadcastindim

Expande as dimensões e/ou classificação de um tensor de entrada, duplicando os dados no tensor operand e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#broadcast_in_dim

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo mlir Descrição
broadcast_dimensions :: mlir :: denseInteLentMattr Atributo de elementos inteiros sem sinais de 64 bits

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor de formato estaticamente do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn tipo f8e4m3fnuz ou tipo f8e4m3b11fnUz ou f8e4m2 flutuação ou de 32 bits ou 64 bits Float ou BFLOAT16 ou Pred (também conhecido com elementos flutuantes de 32 bits ou flutuação de 64 bits ou 2/4/8/8/16/32 bits de uniformes, inteiro signo quantizado ou 2/4/8/8/16/32 bits quantizados inteiros não assinados ou 2/4/8/8/ Uniforme de 16/32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.case (MHLO :: Caseop)

Operação de caso

Produz a saída da execução exatamente uma function de branches dependendo do valor do index .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#case

Exemplo:

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

Características: RecursiveMemoryEffects , SingleBlockImplicitTerminator<ReturnOp> , SingleBlock

Interfaces: InferTypeOpInterface

Operando:

Operando Descrição
index Tensor de valores inteiros sem sinais de 32 bits

Resultados:

Resultado Descrição
"Sem nome" Variada do tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou tipo f8e4m3fn ou tipo f8e4m3fnu ou tipo f8e4m3b11fnUz ou f8e4m3fzz ou tipo f8e4m3b1fnUz ou f8e4m3 Flutuação de 16 bits ou bóia de 32 bits ou 64- bit bit ou bfloat16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou número inteiro ou complexo 2/4/8/16/32/64 bits Digite com elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/8/16/32 bits de uniforme de uniforme ou número inteiro signo ou 2/4/8/16/32 bits quantizados valores inteiros não assinados ou tensor classificado de 2 /4/8/8/16/32 bits quantizados por eixo inteiro assinado ou 2/4/8/16/32 bits quantizados por eixo valores inteiros não assinados ou token

mhlo.cbrt (MHLO :: CBRTOP)

Operação CBRT

Sintaxe:

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

Executa a operação de raiz cúbica em termos de elemento no tensor operand e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cbrt

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType opesondSandResulttype, Elementwise , SameOperandsAndResultShape

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

Efeitos: MemoryEffects::Effect{}

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou tipo BFLOAT16 ou tipo de complexo com elementos flutuantes de 32 bits ou de 64 bits ou 2/4/8/16/32 bits de uniforme, número inteiro signo quantizado ou 2/4/8/16/32 bits quantizados valores inteiros quantizados não assinados

Resultados:

Resultado Descrição
result Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou tipo BFLOAT16 ou tipo de complexo com elementos flutuantes de 32 bits ou de 64 bits ou 2/4/8/16/32 bits de uniforme, número inteiro signo quantizado ou 2/4/8/16/32 bits quantizados valores inteiros quantizados não assinados

mhlo.ceil (MHLO :: CEILOP)

Operação do teto

Sintaxe:

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

Executa tensor de operand no elemento e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#ceil

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType opesondSandResulttype, Elementwise , SameOperandsAndResultShape

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

Efeitos: MemoryEffects::Effect{}

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 ou 2/4/8/16/32 bits de uniforme, número inteiro signo quantizado ou 2/4/8/16/32 bits de uniforme quantizado valores inteiros não assinados quantizados

Resultados:

Resultado Descrição
result Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 ou 2/4/8/16/32 bits de uniforme, número inteiro signo quantizado ou 2/4/8/16/32 bits de uniforme quantizado valores inteiros não assinados quantizados

mhlo.cholesky (MHLO :: Choleskyop)

Operação Cholesky

Calcula a decomposição de Cholesky de um lote de matrizes.

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#cholesky

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , InferTensorType , SameOperandsAndResultElementType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo mlir Descrição
lower :: mlir :: boolattr atributo bool

Operando:

Operando Descrição
a Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou tipo bfloat16 ou tipo complexo com flutuação de 32 bits ou elementos de flutuação de 64 bits valores

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou tipo bfloat16 ou tipo complexo com flutuação de 32 bits ou elementos de flutuação de 64 bits valores

mhlo.clamp (MHLO :: Clampop)

Operação de grampo

Sintaxe:

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

Prenda todos os elementos do tensor de operand entre um valor mínimo e máximo e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#clamp

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , HLO_BroadcastingElementwise , InferTensorType , SameOperandsAndResultElementType

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

Efeitos: MemoryEffects::Effect{}

Operando:

Operando Descrição
min Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados
max Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
result Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.collective_broadcast (MHLO :: CollectiveBroadCastop)

Operação CollectiveBroadcast

Dentro de cada grupo de processo na grade do processo, envie o valor do tensor do operand do processo de origem para os processos de destino e produza um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_broadcast

Exemplo:

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

Características: CompatibleOperandsAndResultType

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Atributos:

Atributo Tipo mlir Descrição
replica_groups :: mlir :: denseInteLentMattr Atributo de elementos inteiros sem sinais de 64 bits
channel_handle :: mlir :: mhlo :: ChannelHandleattr Dois números inteiros de 64 bits 'Handle' e 'Type'

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.collective_permute (mhlo :: collectivepermuteop)

Operação CollectivePermute

Dentro de cada grupo de processo na grade do processo, envia o valor do tensor do operand do processo de origem para o processo de destino e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#collective_permute

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait especulável CompatibleOperandsAndResultType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo mlir Descrição
source_target_pairs :: mlir :: denseInteLentMattr Atributo de elementos inteiros sem sinais de 64 bits
channel_handle :: mlir :: mhlo :: ChannelHandleattr Dois números inteiros de 64 bits 'Handle' e 'Type'

Operando:

Operando Descrição
operand Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

mhlo.compare (MHLO :: Compareop)

Compare operação

Sintaxe:

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

Executa a comparação de elementos de tensores de lhs e rhs de acordo com comparison_direction e compare_type , e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#compare

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , Elementwise , InferTensorType , SameOperandsAndResultShape , SameOperandsElementType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo mlir Descrição
comparison_direction :: mlir :: mhlO :: ComparisOndirectionAttr Qual operação de comparação para executar.
compare_type :: mlir :: mhlo :: comparonTypeattr Qual tipo de comparação a ser usado.

Operando:

Operando Descrição
lhs Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados
rhs Tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou f8e4m3fn ou tipo f8e4m3fnUz ou f8e4b1fnUz ou f8e5m2 tipo ou 4 ou flutuação de 32 bits ou flutuação de 64 bits ou BFLOAT16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bit Elementos flutuantes de bóia de 32 bits ou de 64 bits ou 2/4/8/16/32 bits quantizados inteiros assinados quantizados ou 2/4/8/16/32 bits quantizados inteiros não assinados quantizados ou 2/4/8/16 /Uniforme de 32 bits quantizado por eixo inteiro assinado ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" Tensor classificado de valores pred (também conhecidos como boolean ou 1 bit)

mhlo.complex (MHLO :: Complexop)

Operação complexa

Sintaxe:

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

Executa conversão de elemento para um valor complexo de um par de valores reais e imaginários, lhs e rhs e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#complex

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape , SameOperandsElementType

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

Efeitos: MemoryEffects::Effect{}

Operando:

Operando Descrição
lhs Tensor classificado de flutuação de 32 bits ou valores de flutuação de 64 bits
rhs Tensor classificado de flutuação de 32 bits ou valores de flutuação de 64 bits

Resultados:

Resultado Descrição
result Tensor classificado do tipo complexo com flutuação de 32 bits ou elementos de flutuação de 64 bits valores

mhlo.composite (MHLO :: CompositeOp)

Operação composta

Sintaxe:

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

Encapsula uma operação composta (composta) de outras operações da StableHlo, recebendo inputs e composite_attributes e produzindo results . A semântica do OP é implementada pelo atributo decomposition . O OP composite pode ser substituído por sua decomposição sem alterar a semântica do programa. Nos casos em que a decomposição não fornece a mesma semântica do OP, prefira usar custom_call .

O campo version (padrão para 0 ) é usado para denotar quando a semântica de um composto muda.

Consulte: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#composite

Exemplo:

%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

Atributos:

Atributo Tipo mlir Descrição
name :: mlir :: stringattr atributo da string
composite_attributes :: mlir :: dictionyAttr Dicionário de valores de atributo nomeado
decomposition :: mlir :: flatsymbolRefattr atributo de referência de símbolo plano
version :: mlir :: integerattr Atributo inteiro sem sinais de 32 bits

Operando:

Operando Descrição
inputs Variada do tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou tipo f8e4m3fn ou tipo f8e4m3fnu ou tipo f8e4m3b11fnUz ou f8e4m3fzz ou tipo f8e4m3b1fnUz ou f8e4m3 Flutuação de 16 bits ou bóia de 32 bits ou 64- bit bit ou bfloat16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou número inteiro ou complexo 2/4/8/16/32/64 bits Digite com elementos de flutuação de 32 bits ou de 64 bits ou 2/4/8/8/16/32 bits de uniforme, número inteiro assinado ou 2/4/8/16/32 bits quantizados não assinados ou 2/4/8 O uniforme/16/32 bits quantizado por eixo assinado inteiro ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados ou token ou tupla aninhada com qualquer combinação de tensor classificado do tipo f4e2m1fn ou f6e2m3fn ou f6e3m2fn type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1 -Bit inteiro) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bits não assinados ou tipo complexo com flutuação de 32 bits ou elementos de bóia de 64 bits ou 2/4/8/16/32 bits uniformizados inteiros sinalizados quantizados ou 2/4/8/8/16/32 bits quantizados valores inteiros quantizados quantizados ou tensor de 2/4/8/32 de bits por eixo inteiro assinado ou 2/4/8/16/32 bits quantizados por eixo valores inteiros não assinados ou valores de token

Resultados:

Resultado Descrição
"Sem nome" Variada do tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou tipo f8e4m3fn ou tipo f8e4m3fnu ou tipo f8e4m3b11fnUz ou f8e4m3fzz ou tipo f8e4m3b1fnUz ou f8e4m3 Flutuação de 16 bits ou bóia de 32 bits ou 64- bit bit ou bfloat16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou número inteiro ou complexo 2/4/8/16/32/64 bits Digite com elementos de flutuação de 32 bits ou de 64 bits ou 2/4/8/8/16/32 bits de uniforme, número inteiro assinado ou 2/4/8/16/32 bits quantizados não assinados ou 2/4/8 O uniforme/16/32 bits quantizado por eixo assinado inteiro ou uniforme de 2/4/8/16/32 bits quantizado por eixo valores inteiros não assinados ou token ou tupla aninhada com qualquer combinação de tensor classificado do tipo f4e2m1fn ou f6e2m3fn ou f6e3m2fn type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1 -Bit inteiro) ou 2/4/8/16/32/64 bits sem sinais ou 2/4/8/16/32/64 bits não assinados ou tipo complexo com flutuação de 32 bits ou elementos de bóia de 64 bits ou 2/4/8/16/32 bits uniformizados inteiros sinalizados quantizados ou 2/4/8/8/16/32 bits quantizados valores inteiros quantizados quantizados ou tensor de 2/4/8/32 de bits por eixo inteiro assinado ou 2/4/8/16/32 bits quantizados por eixo valores inteiros não assinados ou valores de token

mhlo.concatenate (MHLO :: Concatenateop)

Operação concatenada

Concatena um número variádico de tensores em inputs ao longo da dimensão dimension na mesma ordem que os argumentos fornecidos e produz um tensor result .

Veja: https://github.com/openxla/stablehlo/blob/main/docs/spec.md#concatenate

Exemplo:

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

Traços: AlwaysSpeculatableImplTrait , SameOperandsAndResultElementType

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

Efeitos: MemoryEffects::Effect{}

Atributos:

Atributo Tipo mlir Descrição
dimension :: mlir :: integerattr Atributo inteiro sem sinais de 64 bits cujo valor não é negativo

Operando:

Operando Descrição
val Variada do tensor classificado do tipo F4e2m1fn ou tipo f6e2m3fn ou tipo f6e3m2fn ou tipo f8e3m4 ou tipo f8e4m3 ou tipo f8e4m3fn ou tipo f8e4m3fnu ou tipo f8e4m3b11fnUz ou f8e4m3fzz ou tipo f8e4m3b1fnUz ou f8e4m3 Flutuação de 16 bits ou bóia de 32 bits ou 64- bit bit ou bfloat16 tipo ou pred (também conhecido como booleano ou número inteiro de 1 bit) ou 2/4/8/16/32/64 bits sem sinais ou número inteiro ou complexo 2/4/8/16/32/64 bits Digite com elementos de flutuação de 32 bits ou de 64 bits ou 2/4/8/8/16/32 bits de uniforme, número inteiro assinado ou 2/4/8/16/32 bits quantizados não assinados ou 2/4/8 /16/32 bits quantizados por eixo inteiro assinado ou 2/4/8/16/32 bits quantizados por eixo valores inteiros não assinados

Resultados:

Resultado Descrição
"Sem nome" ranked tensor of f4E2M1FN type or f6E2M3FN type or f6E3M2FN type or f8E3M4 type or f8E4M3 type or f8E4M3FN type or f8E4M3FNUZ type or f8E4M3B11FNUZ type or f8E5M2 type or f8E5M2FNUZ type or f8E8M0FNU type or 16-bit float or 32-bit float or 64-bit float or bfloat16 type or pred (AKA boolean or 1-bit integer) or 2/4/8/16/32/64-bit signless integer or 2/4/8/16/32/64-bit unsigned integer or complex type with 32-bit float or 64-bit float elements or 2/4/8/16/32-bit uniform quantized signed integer or 2/4/8/16/32-bit uniform quantized unsigned integer or 2/4/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

Exemplo:

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

Atributo MLIR Type Descrição
value ::mlir::ElementsAttr constant vector/tensor attribute

Results:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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.

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

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

Effects: MemoryEffects::Effect{}

Attributes:

Atributo MLIR Type Descrição
cross_program_prefetch_index ::mlir::IntegerAttr 32-bit signless integer attribute

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%output = mhlo.create_token : !mhlo.token

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:

Resultado Descrição
output ficha

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

Exemplo:

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

Atributo MLIR Type Descrição
replica_groups ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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:

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Atributo MLIR Type Descrição
precision_config ::mlir::ArrayAttr Precision Config attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
dimension_numbers ::mlir::mhlo::GatherDimensionNumbersAttr Attribute that models the dimension information for gather
indices_are_sorted ::mlir::BoolAttr bool attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferShapedTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Atributo MLIR Type Descrição
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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 Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%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 Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
slice_sizes ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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

Results:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%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 Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
einsum_config ::mlir::StringAttr string attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
fft_type ::mlir::mhlo::FftTypeAttr XLA fast fourier transform type.
fft_length ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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:

Atributo MLIR Type Descrição
fusion_kind ::mlir::mhlo::FusionKindAttr fusion kind
output_operand_aliases ::mlir::ArrayAttr Aliasing attribute for outputs and operands of Fusion

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Operand Descrição
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:

Resultado Descrição
«unnamed» tensor of 32-bit signless integer values

mhlo.get_tuple_element (mhlo::GetTupleElementOp)

GetTupleElement operation

Sintaxe:

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

Exemplo:

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

Atributo MLIR Type Descrição
index ::mlir::IntegerAttr 32-bit signless integer attribute whose value is non-negative

Operands:

Operand Descrição
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:

Resultado Descrição
«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:

Operand Descrição
pred ranked tensor of pred (AKA boolean or 1-bit integer) values

Results:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Attributes:

Atributo MLIR Type Descrição
infeed_config ::mlir::StringAttr string attribute
layout ::mlir::ArrayAttr array attribute

Operands:

Operand Descrição
token ficha

Results:

Resultado Descrição
«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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Attributes:

Atributo MLIR Type Descrição
iota_dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Results:

Resultado Descrição
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 valores

mhlo.is_finite (mhlo::IsFiniteOp)

IsFinite operation

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
y ranked tensor of pred (AKA boolean or 1-bit integer) values

mhlo.log (mhlo::LogOp)

Log operation

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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:

Operand Descrição
shapes variadic of 1D tensor of index values

Results:

Resultado Descrição
results variadic of 1D tensor of index values

mhlo.multiply (mhlo::MulOp)

Mul operation

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%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 Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Interfaces: InferTypeOpInterface

Attributes:

Atributo MLIR Type Descrição
outfeed_config ::mlir::StringAttr string attribute

Operands:

Operand Descrição
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 ficha

Results:

Resultado Descrição
«unnamed» ficha

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

Exemplo:

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

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

%result = mhlo.partition_id : tensor<ui32>

Interfaces: InferTypeOpInterface

Results:

Resultado Descrição
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.popcnt (mhlo::PopulationCountOp)

PopulationCount operation

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%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 Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

Operand Descrição
token ficha

Results:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise

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

Effects: MemoryEffects::Effect{}

Attributes:

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%result = mhlo.replica_id : tensor<ui32>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Results:

Resultado Descrição
«unnamed» ranked tensor of 32-bit unsigned integer values

mhlo.reshape (mhlo::ReshapeOp)

Reshape operation

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , HLO_CompatibleOperandsAndResultElementType

Interfaces: ConditionallySpeculatable , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
«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:

Atributo MLIR Type Descrição
dimensions ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Traits: InferTensorType

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Atributo MLIR Type Descrição
rng_distribution ::mlir::mhlo::RngDistributionAttr XLA PRNG distribution to be used.

Operands:

Operand Descrição
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

Results:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
rng_algorithm ::mlir::mhlo::RngAlgorithmAttr XLA PRNG algorithm to be used.

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
channel_handle ::mlir::mhlo::ChannelHandleAttr two 64-bit integers 'handle' and 'type'
is_host_transfer ::mlir::BoolAttr bool attribute

Operands:

Operand Descrição
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 ficha

Results:

Resultado Descrição
«unnamed» ficha

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

Exemplo:

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

Atributo MLIR Type Descrição
dimension ::mlir::IntegerAttr 64-bit signless integer attribute whose value is non-negative

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
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 Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
dimension ::mlir::IntegerAttr 64-bit signless integer attribute
is_stable ::mlir::BoolAttr bool attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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:

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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 Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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.

Exemplo:

%0 = mhlo.tan %arg0 : tensor<2xf32>

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%values, %indices = mhlo.topk(%operand, k=5, largest=true)
  : tensor<100xf32> -> (tensor<5xf32>, tensor<5xi32>)

Traits: InferTensorType , RecursiveMemoryEffects

Interfaces: InferShapedTypeOpInterface , InferTypeOpInterface

Attributes:

Atributo MLIR Type Descrição
k ::mlir::IntegerAttr 64-bit signless integer attribute
largest ::mlir::BoolAttr bool attribute

Operands:

Operand Descrição
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:

Resultado Descrição
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.

Exemplo:

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

Atributo MLIR Type Descrição
dim ::mlir::IntegerAttr 64-bit signless integer attribute
batch_dims ::mlir::IntegerAttr 64-bit signless integer attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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.

Exemplo:

mhlo.trace %arg0, "In test code." : tensor<5x1x5xi32>

Attributes:

Atributo MLIR Type Descrição
tag ::mlir::StringAttr string attribute

Operands:

Operand Descrição
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

Exemplo:

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

Atributo MLIR Type Descrição
permutation ::mlir::DenseIntElementsAttr 64-bit signless integer elements attribute

Operands:

Operand Descrição
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:

Resultado Descrição
«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

Exemplo:

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

Atributo MLIR Type Descrição
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:

Operand Descrição
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:

Resultado Descrição
«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

Sintaxe:

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

Exemplo:

%result = mhlo.tuple %val0, %val1 : tuple<tensor<2xf32>, tuple<tensor<i32>>>

Traits: AlwaysSpeculatableImplTrait

Interfaces: ConditionallySpeculatable , InferTypeOpInterface , NoMemoryEffect (MemoryEffectOpInterface)

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%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 Descrição
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:

Resultado Descrição
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

Sintaxe:

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

Exemplo:

%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 Descrição
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

Results:

Resultado Descrição
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

Exemplo:

%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 Descrição
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:

Resultado Descrição
«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

Sintaxe:

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:

Atributo MLIR Type Descrição
delta ::mlir::IntegerAttr 64-bit signless integer attribute

Results:

Resultado Descrição
«unnamed» statically shaped tensor of 64-bit unsigned integer values

mhlo.xor (mhlo::XorOp)

Xor operation

Sintaxe:

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

Exemplo:

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

Traits: AlwaysSpeculatableImplTrait , Commutative , CompatibleOperandsAndResultType , Elementwise , SameOperandsAndResultShape

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

Effects: MemoryEffects::Effect{}

Operands:

Operand Descrição
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:

Resultado Descrição
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

Atributos

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 ...
}

Parâmetros:

Parâmetro C++ type Descrição
argTupleIndices ::llvm::ArrayRef<int64_t> Dimensão
resultIndex int64_t
resultTupleIndices ::llvm::ArrayRef<int64_t> Dimensão
isMustAlias bool

ChannelHandleAttr

two 64-bit integers 'handle' and 'type'

Sintaxe:

#mhlo.channel_handle<
  int64_t,   # handle
  int64_t   # type
>

Parâmetros:

Parâmetro C++ type Descrição
lidar int64_t
tipo int64_t

ComparisonDirectionAttr

Which comparison operation to perform.

Sintaxe:

#mhlo.comparison_direction<
  ::mlir::mhlo::ComparisonDirection   # value
>

Enum cases:

  • EQ ( EQ )
  • NE ( NE )
  • GE ( GE )
  • GT ( GT )
  • LE ( LE )
  • LT ( LT )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::ComparisonDirection an enum of type ComparisonDirection

ComparisonTypeAttr

Which comparison type to use.

Sintaxe:

#mhlo.comparison_type<
  ::mlir::mhlo::ComparisonType   # value
>

Enum cases:

  • NOTYPE ( NOTYPE )
  • FLOAT ( FLOAT )
  • TOTALORDER ( TOTALORDER )
  • SIGNED ( SIGNED )
  • UNSIGNED ( UNSIGNED )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::ComparisonType an enum of type ComparisonType

ConvDimensionNumbersAttr

Structure of dimension information for conv op

Parâmetros:

Parâmetro C++ type Descrição
inputBatchDimension int64_t
inputFeatureDimension int64_t
inputSpatialDimensions ::llvm::ArrayRef<int64_t> Dimensão
kernelInputFeatureDimension int64_t
kernelOutputFeatureDimension int64_t
kernelSpatialDimensions ::llvm::ArrayRef<int64_t> Dimensão
outputBatchDimension int64_t
outputFeatureDimension int64_t
outputSpatialDimensions ::llvm::ArrayRef<int64_t> Dimensão

CrossProgramPrefetchAttr

Argument that is prefetched from another program

Sintaxe:

#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.

Por exemplo,

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.

Parâmetros:

Parâmetro C++ type Descrição
parâmetro int64_t
índices ::llvm::ArrayRef<int64_t> Dimensão
desvio std::optional<int64_t>

CustomCallScheduleAttr

Specifies the desired schedule for the custom-call.

Sintaxe:

#mhlo.custom_call_schedule<
  ::mlir::mhlo::CustomCallSchedule   # value
>

Enum cases:

  • NONE ( NONE )
  • LATEST ( LATEST )
  • EARLIEST ( EARLIEST )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::CustomCallSchedule an enum of type CustomCallSchedule

DequantizeModeAttr

Dequantization mode. Only MIN_COMBINED is supported.

Sintaxe:

#mhlo.dequantize_mode<
  ::mlir::mhlo::DequantizeMode   # value
>

Enum cases:

  • MIN_COMBINED ( MIN_COMBINED )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::DequantizeMode an enum of type DequantizeMode

DomainKindAttr

Kind of domain metatdata attached to an HLO domain.

Sintaxe:

#mhlo.kind<
  ::mlir::mhlo::DomainKind   # value
>

Enum cases:

  • sharding ( sharding )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::DomainKind an enum of type DomainKind

DotAlgorithmAttr

Attribute that models the algorithm constraints to use for computing dot.

Sintaxe:

#mhlo.dot_algorithm<
  Type,   # lhsPrecisionType
  Type,   # rhsPrecisionType
  Type,   # accumulationType
  int64_t,   # lhsComponentCount
  int64_t,   # rhsComponentCount
  int64_t,   # numPrimitiveOperations
  bool   # allowImpreciseAccumulation
>

Parâmetros:

Parâmetro C++ type Descrição
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.

Parâmetros:

Parâmetro C++ type Descrição
lhsBatchingDimensions ::llvm::ArrayRef<int64_t> Dimensão
rhsBatchingDimensions ::llvm::ArrayRef<int64_t> Dimensão
lhsContractingDimensions ::llvm::ArrayRef<int64_t> Dimensão
rhsContractingDimensions ::llvm::ArrayRef<int64_t> Dimensão

FftTypeAttr

XLA fast fourier transform type.

Sintaxe:

#mhlo.fft_type<
  ::mlir::mhlo::FftType   # value
>

Enum cases:

  • FFT ( FFT )
  • IFFT ( IFFT )
  • RFFT ( RFFT )
  • IRFFT ( IRFFT )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::FftType an enum of type FftType

FusionKindAttr

fusion kind

Sintaxe:

#mhlo.fusion_kind<
  ::mlir::mhlo::FusionKind   # value
>

Enum cases:

  • kLoop ( kLoop )
  • kInput ( kInput )
  • kOutput ( kOutput )
  • kCustom ( kCustom )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::FusionKind an enum of type FusionKind

GatherDimensionNumbersAttr

Attribute that models the dimension information for gather

Parâmetros:

Parâmetro C++ type Descrição
offsetDims ::llvm::ArrayRef<int64_t> Dimensão
collapsedSliceDims ::llvm::ArrayRef<int64_t> Dimensão
operandBatchingDims ::llvm::ArrayRef<int64_t> Dimensão
startIndicesBatchingDims ::llvm::ArrayRef<int64_t> Dimensão
startIndexMap ::llvm::ArrayRef<int64_t> Dimensão
indexVectorDim int64_t

OutputOperandAliasAttr

Attribute that models the alias relationship of output and operand of a CustomCall op

Sintaxe:

#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>.

Parâmetros:

Parâmetro C++ type Descrição
outputTupleIndices ::llvm::ArrayRef<int64_t> Dimensão
operandIndex int64_t
operandTupleIndices ::llvm::ArrayRef<int64_t> Dimensão

PrecisionAttr

XLA precision for an operand. Has backend specific meaning.

Sintaxe:

#mhlo.precision<
  ::mlir::mhlo::Precision   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • HIGH ( HIGH )
  • HIGHEST ( HIGHEST )
  • PACKED_NIBBLE ( PACKED_NIBBLE )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::Precision an enum of type Precision

RngAlgorithmAttr

XLA PRNG algorithm to be used.

Sintaxe:

#mhlo.rng_algorithm<
  ::mlir::mhlo::RngAlgorithm   # value
>

Enum cases:

  • DEFAULT ( DEFAULT )
  • THREE_FRY ( THREE_FRY )
  • PHILOX ( PHILOX )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::RngAlgorithm an enum of type RngAlgorithm

RngDistributionAttr

XLA PRNG distribution to be used.

Sintaxe:

#mhlo.rng_distribution<
  ::mlir::mhlo::RngDistribution   # value
>

Enum cases:

  • UNIFORM ( UNIFORM )
  • NORMAL ( NORMAL )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::RngDistribution an enum of type RngDistribution

ScatterDimensionNumbersAttr

Attribute that models the dimension information for scatter

Parâmetros:

Parâmetro C++ type Descrição
updateWindowDims ::llvm::ArrayRef<int64_t> Dimensão
insertedWindowDims ::llvm::ArrayRef<int64_t> Dimensão
inputBatchingDims ::llvm::ArrayRef<int64_t> Dimensão
scatterIndicesBatchingDims ::llvm::ArrayRef<int64_t> Dimensão
scatterDimsToOperandDims ::llvm::ArrayRef<int64_t> Dimensão
indexVectorDim int64_t

SparsityDescriptorAttr

Describes structured (N:M) sparsity configuration

Sintaxe:

#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.

Parâmetros:

Parâmetro C++ type Descrição
dimensão int64_t
n int64_t
eu int64_t

TransposeAttr

Transpose options

Sintaxe:

#mhlo.transpose<
  ::mlir::mhlo::Transpose   # value
>

Enum cases:

  • TRANSPOSE_INVALID ( TRANSPOSE_INVALID )
  • NO_TRANSPOSE ( NO_TRANSPOSE )
  • TRANSPOSE ( TRANSPOSE )
  • ADJOINT ( ADJOINT )

Parâmetros:

Parâmetro C++ type Descrição
valor ::mlir::mhlo::Transpose an enum of type Transpose

TypeExtensionsAttr

Attribute that extends tensor type with MHLO type properties.

Sintaxe:

#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 .

Parâmetros:

Parâmetro C++ type Descrição
limites ::llvm::ArrayRef<int64_t>

Tipos

AsyncBundleType

Opaque collection of other types

Sintaxe:

!mhlo.async_bundle<
  ::llvm::ArrayRef<Type>   # types
>

Parâmetros:

Parâmetro C++ type Descrição
tipos ::llvm::ArrayRef<Type>

Enums

ComparisonDirection

Which comparison operation to perform.

Cases:

Símbolo Valor Corda
equalização 0 equalização
NE 1 NE
GE 2 GE
GT 3 GT
LE 4 LE
LT 5 LT

ComparisonType

Which comparison type to use.

Cases:

Símbolo Valor Corda
NOTYPE 0 NOTYPE
FLUTUADOR 1 FLUTUADOR
TOTALORDER 2 TOTALORDER
SIGNED 3 SIGNED
UNSIGNED 4 UNSIGNED

CustomCallApiVersion

Custom call API version

Cases:

Símbolo Valor Corda
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.

Cases:

Símbolo Valor Corda
NENHUM 0 NENHUM
MAIS RECENTE 1 MAIS RECENTE
EARLIEST 2 EARLIEST

DequantizeMode

Dequantization mode. Only MIN_COMBINED is supported.

Cases:

Símbolo Valor Corda
MIN_COMBINED 0 MIN_COMBINED

DomainKind

Kind of domain metatdata attached to an HLO domain.

Cases:

Símbolo Valor Corda
sharding 0 sharding

FftType

XLA fast fourier transform type.

Cases:

Símbolo Valor Corda
FFT 0 FFT
IFFT 1 IFFT
RFFT 2 RFFT
IRFFT 3 IRFFT

FusionKind

fusion kind

Cases:

Símbolo Valor Corda
kLoop 0 kLoop
kInput 1 kInput
kOutput 2 kOutput
kCustom 3 kCustom

Precisão

XLA precision for an operand. Has backend specific meaning.

Cases:

Símbolo Valor Corda
PADRÃO 0 PADRÃO
ALTO 1 ALTO
HIGHEST 2 HIGHEST
PACKED_NIBBLE 3 PACKED_NIBBLE

RngAlgorithm

XLA PRNG algorithm to be used.

Cases:

Símbolo Valor Corda
PADRÃO 0 PADRÃO
THREE_FRY 1 THREE_FRY
PHILOX 2 PHILOX

RngDistribution

XLA PRNG distribution to be used.

Cases:

Símbolo Valor Corda
UNIFORME 1 UNIFORME
NORMAL 2 NORMAL

Transpor

Transpose options

Cases:

Símbolo Valor Corda
TRANSPOSE_INVALID 0 TRANSPOSE_INVALID
NO_TRANSPOSE 1 NO_TRANSPOSE
TRANSPOR 2 TRANSPOR
ADJOINT 3 ADJOINT