Interpreter Error

[Axel Wong]

Hi everyone,

I've encountered a runtime error with the following IR. There is no error during IR module verifying.

module {
  func.func @main(%arg0: tensor<1x3x32x32xf32>, %arg1: tensor<6x3x5x5xf32>, %arg2: tensor<6xf32>, %arg3: tensor<16x6x5x5xf32>, %arg4: tensor<16xf32>, %arg5: tensor<120x400xf32>, %arg6: tensor<120xf32>, %arg7: tensor<84x120xf32>, %arg8: tensor<84xf32>, %arg9: tensor<10x84xf32>, %arg10: tensor<10xf32>) -> tensor<1x10xf32> {
    %0 = stablehlo.uniform_quantize %arg0 : (tensor<1x3x32x32xf32>) -> tensor<1x3x32x32x!quant.uniform<i8:f32, 0.029222087934613228>>
    %1 = stablehlo.uniform_quantize %arg1 : (tensor<6x3x5x5xf32>) -> tensor<6x3x5x5x!quant.uniform<i8:f32, 0.0039700339548289776>>
    %2 = stablehlo.uniform_quantize %arg2 : (tensor<6xf32>) -> tensor<6x!quant.uniform<i8:f32, 0.0021052095107734203>>
    %3 = stablehlo.convolution(%0, %1) dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1], window = {stride = [1, 1], pad = [[0, 0], [0, 0]], lhs_dilate = [1, 1], rhs_dilate = [1, 1]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x3x32x32x!quant.uniform<i8:f32, 0.029222087934613228>>, tensor<6x3x5x5x!quant.uniform<i8:f32, 0.0039700339548289776>>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %4 = stablehlo.broadcast_in_dim %2, dims = [1] : (tensor<6x!quant.uniform<i8:f32, 0.0021052095107734203>>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.0021052095107734203>>
    %5 = stablehlo.add %3, %4 : (tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>, tensor<1x6x28x28x!quant.uniform<i8:f32, 0.0021052095107734203>>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %cst = stablehlo.constant dense<0.000000e+00> : tensor<1x6x28x28xf32>
    %6 = stablehlo.uniform_quantize %cst : (tensor<1x6x28x28xf32>) -> tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %7 = stablehlo.maximum %5, %6 : tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>
    %cst_0 = stablehlo.constant dense<0xFF800000> : tensor<f32>
    %8 = stablehlo.uniform_quantize %cst_0 : (tensor<f32>) -> tensor<!quant.uniform<i8:f32, 0.053985215723514557>>
    %9 = "stablehlo.reduce_window"(%7, %8) <{base_dilations = array<i64: 1, 1, 1, 1>, padding = dense<0> : tensor<4x2xi64>, window_dilations = array<i64: 1, 1, 1, 1>, window_dimensions = array<i64: 1, 1, 2, 2>, window_strides = array<i64: 1, 1, 2, 2>}> ({
    ^bb0(%arg11: tensor<!quant.uniform<i8:f32, 0.053985215723514557>>, %arg12: tensor<!quant.uniform<i8:f32, 0.053985215723514557>>):
      %37 = stablehlo.maximum %arg11, %arg12 : tensor<!quant.uniform<i8:f32, 0.053985215723514557>>
      stablehlo.return %37 : tensor<!quant.uniform<i8:f32, 0.053985215723514557>>
    }) : (tensor<1x6x28x28x!quant.uniform<i8:f32, 0.053985215723514557>>, tensor<!quant.uniform<i8:f32, 0.053985215723514557>>) -> tensor<1x6x14x14x!quant.uniform<i8:f32, 0.04167218878865242>>
    %10 = stablehlo.uniform_quantize %arg3 : (tensor<16x6x5x5xf32>) -> tensor<16x6x5x5x!quant.uniform<i8:f32, 0.016297046095132828>>
    %11 = stablehlo.uniform_quantize %arg4 : (tensor<16xf32>) -> tensor<16x!quant.uniform<i8:f32, 0.004838157445192337>>
    %12 = stablehlo.convolution(%9, %10) dim_numbers = [b, f, 0, 1]x[o, i, 0, 1]->[b, f, 0, 1], window = {stride = [1, 1], pad = [[0, 0], [0, 0]], lhs_dilate = [1, 1], rhs_dilate = [1, 1]} {batch_group_count = 1 : i64, feature_group_count = 1 : i64} : (tensor<1x6x14x14x!quant.uniform<i8:f32, 0.04167218878865242>>, tensor<16x6x5x5x!quant.uniform<i8:f32, 0.016297046095132828>>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %13 = stablehlo.broadcast_in_dim %11, dims = [1] : (tensor<16x!quant.uniform<i8:f32, 0.004838157445192337>>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.004838157445192337>>
    %14 = stablehlo.add %12, %13 : (tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>, tensor<1x16x10x10x!quant.uniform<i8:f32, 0.004838157445192337>>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %cst_1 = stablehlo.constant dense<0.000000e+00> : tensor<1x16x10x10xf32>
    %15 = stablehlo.uniform_quantize %cst_1 : (tensor<1x16x10x10xf32>) -> tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %16 = stablehlo.maximum %14, %15 : tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>
    %cst_2 = stablehlo.constant dense<0xFF800000> : tensor<f32>
    %17 = stablehlo.uniform_quantize %cst_2 : (tensor<f32>) -> tensor<!quant.uniform<i8:f32, 0.20967985689640045>>
    %18 = "stablehlo.reduce_window"(%16, %17) <{base_dilations = array<i64: 1, 1, 1, 1>, padding = dense<0> : tensor<4x2xi64>, window_dilations = array<i64: 1, 1, 1, 1>, window_dimensions = array<i64: 1, 1, 2, 2>, window_strides = array<i64: 1, 1, 2, 2>}> ({
    ^bb0(%arg11: tensor<!quant.uniform<i8:f32, 0.20967985689640045>>, %arg12: tensor<!quant.uniform<i8:f32, 0.20967985689640045>>):
      %37 = stablehlo.maximum %arg11, %arg12 : tensor<!quant.uniform<i8:f32, 0.20967985689640045>>
      stablehlo.return %37 : tensor<!quant.uniform<i8:f32, 0.20967985689640045>>
    }) : (tensor<1x16x10x10x!quant.uniform<i8:f32, 0.20967985689640045>>, tensor<!quant.uniform<i8:f32, 0.20967985689640045>>) -> tensor<1x16x5x5x!quant.uniform<i8:f32, 0.11828774958848953>>
    %19 = stablehlo.reshape %18 : (tensor<1x16x5x5x!quant.uniform<i8:f32, 0.11828774958848953>>) -> tensor<1x400x!quant.uniform<i8:f32, 0.11828774958848953>>
    %20 = stablehlo.uniform_quantize %arg5 : (tensor<120x400xf32>) -> tensor<120x400x!quant.uniform<i8:f32, 0.0099556939676403999>>
    %21 = stablehlo.uniform_quantize %arg6 : (tensor<120xf32>) -> tensor<120x!quant.uniform<i8:f32, 0.0091858189553022385>>
    %22 = stablehlo.transpose %20, dims = [1, 0] : (tensor<120x400x!quant.uniform<i8:f32, 0.0099556939676403999>>) -> tensor<400x120x!quant.uniform<i8:f32, 0.0099556939676403999>>
    %23 = stablehlo.dot_general %19, %22, contracting_dims = [1] x [0] : (tensor<1x400x!quant.uniform<i8:f32, 0.11828774958848953>>, tensor<400x120x!quant.uniform<i8:f32, 0.0099556939676403999>>) -> tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>
    %cst_3 = stablehlo.constant dense<0.000000e+00> : tensor<1x120xf32>
    %24 = stablehlo.uniform_quantize %cst_3 : (tensor<1x120xf32>) -> tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>
    %25 = stablehlo.maximum %23, %24 : tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>
    %26 = stablehlo.uniform_quantize %arg7 : (tensor<84x120xf32>) -> tensor<84x120x!quant.uniform<i8:f32, 0.011763181537389755>>
    %27 = stablehlo.uniform_quantize %arg8 : (tensor<84xf32>) -> tensor<84x!quant.uniform<i8:f32, 0.008067096583545208>>
    %28 = stablehlo.transpose %26, dims = [1, 0] : (tensor<84x120x!quant.uniform<i8:f32, 0.011763181537389755>>) -> tensor<120x84x!quant.uniform<i8:f32, 0.011763181537389755>>
    %29 = stablehlo.dot_general %25, %28, contracting_dims = [1] x [0] : (tensor<1x120x!quant.uniform<i8:f32, 0.46472719311714172>>, tensor<120x84x!quant.uniform<i8:f32, 0.011763181537389755>>) -> tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>
    %cst_4 = stablehlo.constant dense<0.000000e+00> : tensor<1x84xf32>
    %30 = stablehlo.uniform_quantize %cst_4 : (tensor<1x84xf32>) -> tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>
    %31 = stablehlo.maximum %29, %30 : tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>
    %32 = stablehlo.uniform_quantize %arg9 : (tensor<10x84xf32>) -> tensor<10x84x!quant.uniform<i8:f32, 0.013255398720502853>>
    %33 = stablehlo.uniform_quantize %arg10 : (tensor<10xf32>) -> tensor<10x!quant.uniform<i8:f32, 0.0026047099381685257>>
    %34 = stablehlo.transpose %32, dims = [1, 0] : (tensor<10x84x!quant.uniform<i8:f32, 0.013255398720502853>>) -> tensor<84x10x!quant.uniform<i8:f32, 0.013255398720502853>>
    %35 = stablehlo.dot_general %31, %34, contracting_dims = [1] x [0] : (tensor<1x84x!quant.uniform<i8:f32, 0.40483295917510986>>, tensor<84x10x!quant.uniform<i8:f32, 0.013255398720502853>>) -> tensor<1x10x!quant.uniform<i8:f32, 0.18991459906101227>>
    %36 = stablehlo.uniform_dequantize %35 : (tensor<1x10x!quant.uniform<i8:f32, 0.18991459906101227>>) -> tensor<1x10xf32>
    return %36 : tensor<1x10xf32>
  }
}

LLVM ERROR: Unsupported op: %0 = stablehlo.uniform_quantize %arg0 : (tensor<1x3x32x32xf32>) -> tensor<1x3x32x32x!quant.uniform<i8:f32, 0.029222087934613228>>

Aborted (core dumped)

Appreciate any help and ideas.

Thanks

Axel

[Sandeep Dasgupta]

Hello Axel

I have just added an https://github.com/openxla/stablehlo/issues/2373 to address similar cases. The interpreter does not have support for the quantized types yet and that is something we would like to explore further. 

Also, figuring out the solution presented in the post  may not be applicable directly as the pass mhlo-quant-lelaglize-to-int met some issue which we are still exploring. We will get back asap. 

Regards,

Sandeep

--
You received this message because you are subscribed to the Google Groups "OpenXLA Discuss" group.
To unsubscribe from this group and stop receiving emails from it, send an email to openxla-discu...@openxla.org.
To view this discussion on the web visit https://groups.google.com/a/openxla.org/d/msgid/openxla-discuss/d84e84e3-422c-4765-aa38-07d999c5efdan%40openxla.org.
For more options, visit https://groups.google.com/a/openxla.org/d/optout.